[clang] Adding support for iterator in motion clauses. (PR #159112)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 7 00:47:33 PDT 2025
https://github.com/ShashwathiNavada updated https://github.com/llvm/llvm-project/pull/159112
>From 7bec4d102da2e3c99a8a8d83167471637d94f897 Mon Sep 17 00:00:00 2001
From: Shashwathi N <nshashwa at pe31.hpc.amslabs.hpecorp.net>
Date: Tue, 16 Sep 2025 10:23:50 -0500
Subject: [PATCH 1/5] Adding support for iterator in motion clauses
---
clang/docs/OpenMPSupport.rst | 2 +-
clang/include/clang/AST/OpenMPClause.h | 27 ++++++++-----
clang/include/clang/Basic/OpenMPKinds.def | 1 +
clang/include/clang/Sema/SemaOpenMP.h | 4 +-
clang/lib/AST/OpenMPClause.cpp | 26 +++++++++----
clang/lib/Parse/ParseOpenMP.cpp | 14 +++++++
clang/lib/Sema/SemaOpenMP.cpp | 38 ++++++++++++-------
clang/lib/Sema/TreeTransform.h | 36 ++++++++++++++----
clang/lib/Serialization/ASTReader.cpp | 4 ++
clang/lib/Serialization/ASTWriter.cpp | 4 ++
.../target_update_iterator_ast_print.cpp | 16 ++++++++
11 files changed, 130 insertions(+), 42 deletions(-)
create mode 100644 clang/test/OpenMP/target_update_iterator_ast_print.cpp
diff --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst
index 0bc8590815220..a3e4272ceaf2f 100644
--- a/clang/docs/OpenMPSupport.rst
+++ b/clang/docs/OpenMPSupport.rst
@@ -266,7 +266,7 @@ implementation.
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | has_device_addr clause on target construct | :none:`unclaimed` | |
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
-| device | iterators in map clause or motion clauses | :none:`unclaimed` | |
+| device | iterators in map clause or motion clauses | :good:`done` | |
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | indirect clause on declare target directive | :part:`In Progress` | |
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index b2a6d4b9182b0..4e5a86482fbd4 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -7476,7 +7476,7 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
/// Motion-modifiers for the 'to' clause.
OpenMPMotionModifierKind MotionModifiers[NumberOfOMPMotionModifiers] = {
- OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown};
+ OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown};
/// Location of motion-modifiers for the 'to' clause.
SourceLocation MotionModifiersLoc[NumberOfOMPMotionModifiers];
@@ -7548,6 +7548,9 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
MotionModifiersLoc[I] = TLoc;
}
+ void setIteratorModifier(Expr *IteratorModifier) {
+ getTrailingObjects<Expr *>()[2 * varlist_size()] = IteratorModifier;
+ }
/// Set colon location.
void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; }
@@ -7556,7 +7559,7 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
size_t numTrailingObjects(OverloadToken<Expr *>) const {
// There are varlist_size() of expressions, and varlist_size() of
// user-defined mappers.
- return 2 * varlist_size();
+ return 2 * varlist_size() + 1;
}
size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
return getUniqueDeclarationsNum();
@@ -7586,7 +7589,7 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists,
- ArrayRef<Expr *> UDMapperRefs,
+ ArrayRef<Expr *> UDMapperRefs,Expr *IteratorModifier,
ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
NestedNameSpecifierLoc UDMQualifierLoc,
@@ -7611,7 +7614,9 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
"Requested modifier exceeds the total number of modifiers.");
return MotionModifiers[Cnt];
}
-
+ Expr *getIteratorModifier() {
+ return getTrailingObjects<Expr *>()[2 * varlist_size()];
+ }
/// Fetches the motion-modifier location at 'Cnt' index of array of modifiers'
/// locations.
///
@@ -7677,7 +7682,7 @@ class OMPFromClause final
/// Motion-modifiers for the 'from' clause.
OpenMPMotionModifierKind MotionModifiers[NumberOfOMPMotionModifiers] = {
- OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown};
+ OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown};
/// Location of motion-modifiers for the 'from' clause.
SourceLocation MotionModifiersLoc[NumberOfOMPMotionModifiers];
@@ -7738,7 +7743,9 @@ class OMPFromClause final
"Unexpected index to store motion modifier, exceeds array size.");
MotionModifiers[I] = T;
}
-
+ void setIteratorModifier(Expr *IteratorModifier) {
+ getTrailingObjects<Expr *>()[2 * varlist_size()] = IteratorModifier;
+ }
/// Set location for the motion-modifier.
///
/// \param I index for motion-modifier location.
@@ -7757,7 +7764,7 @@ class OMPFromClause final
size_t numTrailingObjects(OverloadToken<Expr *>) const {
// There are varlist_size() of expressions, and varlist_size() of
// user-defined mappers.
- return 2 * varlist_size();
+ return 2 * varlist_size() + 1;
}
size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
return getUniqueDeclarationsNum();
@@ -7787,7 +7794,7 @@ class OMPFromClause final
Create(const ASTContext &C, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists,
- ArrayRef<Expr *> UDMapperRefs,
+ ArrayRef<Expr *> UDMapperRefs, Expr *IteratorExpr,
ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId);
@@ -7811,7 +7818,9 @@ class OMPFromClause final
"Requested modifier exceeds the total number of modifiers.");
return MotionModifiers[Cnt];
}
-
+ Expr *getIteratorModifier() {
+ return getTrailingObjects<Expr *>()[2 * varlist_size()];
+ }
/// Fetches the motion-modifier location at 'Cnt' index of array of modifiers'
/// locations.
///
diff --git a/clang/include/clang/Basic/OpenMPKinds.def b/clang/include/clang/Basic/OpenMPKinds.def
index 79c11b851c557..7944dc6916c94 100644
--- a/clang/include/clang/Basic/OpenMPKinds.def
+++ b/clang/include/clang/Basic/OpenMPKinds.def
@@ -194,6 +194,7 @@ OPENMP_MAP_MODIFIER_KIND(ompx_hold)
// Modifiers for 'to' or 'from' clause.
OPENMP_MOTION_MODIFIER_KIND(mapper)
+OPENMP_MOTION_MODIFIER_KIND(iterator)
OPENMP_MOTION_MODIFIER_KIND(present)
// Static attributes for 'dist_schedule' clause.
diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h
index 23827051ed724..ecfbe3284a202 100644
--- a/clang/include/clang/Sema/SemaOpenMP.h
+++ b/clang/include/clang/Sema/SemaOpenMP.h
@@ -1323,7 +1323,7 @@ class SemaOpenMP : public SemaBase {
OMPClause *
ActOnOpenMPToClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
- CXXScopeSpec &MapperIdScopeSpec,
+ Expr *IteratorModifier, CXXScopeSpec &MapperIdScopeSpec,
DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> UnresolvedMappers = {});
@@ -1331,7 +1331,7 @@ class SemaOpenMP : public SemaBase {
OMPClause *
ActOnOpenMPFromClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
- CXXScopeSpec &MapperIdScopeSpec,
+ Expr *IteratorModifier, CXXScopeSpec &MapperIdScopeSpec,
DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> UnresolvedMappers = {});
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index 69d33019c0952..de352df0f93f9 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -1218,7 +1218,7 @@ OMPToClause *OMPToClause::Create(
const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
- ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ Expr *IteratorModifier, ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) {
OMPMappableExprListSizeTy Sizes;
@@ -1240,7 +1240,7 @@ OMPToClause *OMPToClause::Create(
void *Mem = C.Allocate(
totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
- 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ 2 * Sizes.NumVars + 1, Sizes.NumUniqueDeclarations,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
@@ -1250,6 +1250,7 @@ OMPToClause *OMPToClause::Create(
Clause->setVarRefs(Vars);
Clause->setUDMapperRefs(UDMapperRefs);
Clause->setClauseInfo(Declarations, ComponentLists);
+ Clause->setIteratorModifier(IteratorModifier);
return Clause;
}
@@ -1258,17 +1259,19 @@ OMPToClause *OMPToClause::CreateEmpty(const ASTContext &C,
void *Mem = C.Allocate(
totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
- 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ 2 * Sizes.NumVars + 1, Sizes.NumUniqueDeclarations,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
- return new (Mem) OMPToClause(Sizes);
+ OMPToClause *Clause = new (Mem) OMPToClause(Sizes);
+ Clause->setIteratorModifier(nullptr);
+ return Clause;
}
OMPFromClause *OMPFromClause::Create(
const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
- ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ Expr *IteratorModifier, ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) {
OMPMappableExprListSizeTy Sizes;
@@ -1290,7 +1293,7 @@ OMPFromClause *OMPFromClause::Create(
void *Mem = C.Allocate(
totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
- 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ 2 * Sizes.NumVars + 1, Sizes.NumUniqueDeclarations,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
@@ -1301,6 +1304,7 @@ OMPFromClause *OMPFromClause::Create(
Clause->setVarRefs(Vars);
Clause->setUDMapperRefs(UDMapperRefs);
Clause->setClauseInfo(Declarations, ComponentLists);
+ Clause->setIteratorModifier(IteratorModifier);
return Clause;
}
@@ -1310,10 +1314,12 @@ OMPFromClause::CreateEmpty(const ASTContext &C,
void *Mem = C.Allocate(
totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
- 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
+ 2 * Sizes.NumVars + 1, Sizes.NumUniqueDeclarations,
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
- return new (Mem) OMPFromClause(Sizes);
+ OMPFromClause *Clause = new (Mem) OMPFromClause(Sizes);
+ Clause->setIteratorModifier(nullptr);
+ return Clause;
}
void OMPUseDevicePtrClause::setPrivateCopies(ArrayRef<Expr *> VL) {
@@ -2564,12 +2570,16 @@ template <typename T> void OMPClausePrinter::VisitOMPMotionClause(T *Node) {
OS << '(';
for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
if (Node->getMotionModifier(I) != OMPC_MOTION_MODIFIER_unknown) {
+ if (Node->getMotionModifier(I) == OMPC_MOTION_MODIFIER_iterator)
+ PrintIterator(OS, Node, Policy);
+ else {
OS << getOpenMPSimpleClauseTypeName(Node->getClauseKind(),
Node->getMotionModifier(I));
if (Node->getMotionModifier(I) == OMPC_MOTION_MODIFIER_mapper)
PrintMapper(OS, Node, Policy);
if (I < ModifierCount - 1)
OS << ", ";
+ }
}
}
OS << ':';
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 7dceb2d208352..75fd5a12ff3a1 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -4699,6 +4699,19 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
break;
Data.MotionModifiers.push_back(Modifier);
Data.MotionModifiersLoc.push_back(Tok.getLocation());
+ if (PP.getSpelling(Tok) == "iterator" && getLangOpts().OpenMP >= 51) {
+ ColonProtectionRAIIObject ColonRAII(*this);
+ TentativeParsingAction TPA(*this);
+ ExprResult Tail;
+ HasIterator = true;
+ EnterScope(Scope::OpenMPDirectiveScope | Scope::DeclScope);
+ Tail = ParseOpenMPIteratorsExpr();
+ Tail = Actions.ActOnFinishFullExpr(Tail.get(), T.getOpenLocation(),
+ /*DiscardedValue=*/false);
+ if (Tail.isUsable()) {
+ Data.IteratorExpr = Tail.get();
+ }
+ } else {
ConsumeToken();
if (Modifier == OMPC_MOTION_MODIFIER_mapper) {
IsInvalidMapperModifier = parseMapperModifier(Data);
@@ -4712,6 +4725,7 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
// TODO: Is that intentional?
if (Tok.is(tok::comma))
ConsumeToken();
+ }
}
if (!Data.MotionModifiers.empty() && Tok.isNot(tok::colon)) {
if (!IsInvalidMapperModifier) {
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 66bdd6e72d827..f44e6be061c5d 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -17723,16 +17723,16 @@ OMPClause *SemaOpenMP::ActOnOpenMPVarListClause(OpenMPClauseKind Kind,
ExtraModifierLoc, ColonLoc, VarList, Locs);
break;
case OMPC_to:
- Res =
- ActOnOpenMPToClause(Data.MotionModifiers, Data.MotionModifiersLoc,
- Data.ReductionOrMapperIdScopeSpec,
- Data.ReductionOrMapperId, ColonLoc, VarList, Locs);
+ Res = ActOnOpenMPToClause(
+ Data.MotionModifiers, Data.MotionModifiersLoc, Data.IteratorExpr,
+ Data.ReductionOrMapperIdScopeSpec, Data.ReductionOrMapperId, ColonLoc,
+ VarList, Locs);
break;
case OMPC_from:
- Res = ActOnOpenMPFromClause(Data.MotionModifiers, Data.MotionModifiersLoc,
- Data.ReductionOrMapperIdScopeSpec,
- Data.ReductionOrMapperId, ColonLoc, VarList,
- Locs);
+ Res = ActOnOpenMPFromClause(
+ Data.MotionModifiers, Data.MotionModifiersLoc, Data.IteratorExpr,
+ Data.ReductionOrMapperIdScopeSpec, Data.ReductionOrMapperId, ColonLoc,
+ VarList, Locs);
break;
case OMPC_use_device_ptr:
Res = ActOnOpenMPUseDevicePtrClause(VarList, Locs);
@@ -23465,10 +23465,12 @@ void SemaOpenMP::ActOnOpenMPDeclareTargetInitializer(Decl *TargetDecl) {
OMPClause *SemaOpenMP::ActOnOpenMPToClause(
ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
+ Expr *IteratorExpr,
CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
OpenMPMotionModifierKind Modifiers[] = {OMPC_MOTION_MODIFIER_unknown,
+ OMPC_MOTION_MODIFIER_unknown,
OMPC_MOTION_MODIFIER_unknown};
SourceLocation ModifiersLoc[NumberOfOMPMotionModifiers];
@@ -23492,20 +23494,24 @@ OMPClause *SemaOpenMP::ActOnOpenMPToClause(
MapperIdScopeSpec, MapperId, UnresolvedMappers);
if (MVLI.ProcessedVarList.empty())
return nullptr;
-
+ if(IteratorExpr)
+ if (auto *DRE = dyn_cast<DeclRefExpr>(IteratorExpr))
+ if (auto *VD = dyn_cast<VarDecl>(DRE->getDecl()))
+ DSAStack->addIteratorVarDecl(VD);
return OMPToClause::Create(
getASTContext(), Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
- MVLI.VarComponents, MVLI.UDMapperList, Modifiers, ModifiersLoc,
+ MVLI.VarComponents, MVLI.UDMapperList,IteratorExpr, Modifiers, ModifiersLoc,
MapperIdScopeSpec.getWithLocInContext(getASTContext()), MapperId);
}
OMPClause *SemaOpenMP::ActOnOpenMPFromClause(
ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
- ArrayRef<SourceLocation> MotionModifiersLoc,
+ ArrayRef<SourceLocation> MotionModifiersLoc, Expr *IteratorExpr,
CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
OpenMPMotionModifierKind Modifiers[] = {OMPC_MOTION_MODIFIER_unknown,
+ OMPC_MOTION_MODIFIER_unknown,
OMPC_MOTION_MODIFIER_unknown};
SourceLocation ModifiersLoc[NumberOfOMPMotionModifiers];
@@ -23529,11 +23535,15 @@ OMPClause *SemaOpenMP::ActOnOpenMPFromClause(
MapperIdScopeSpec, MapperId, UnresolvedMappers);
if (MVLI.ProcessedVarList.empty())
return nullptr;
-
+ if (IteratorExpr)
+ if (auto *DRE = dyn_cast<DeclRefExpr>(IteratorExpr))
+ if (auto *VD = dyn_cast<VarDecl>(DRE->getDecl()))
+ DSAStack->addIteratorVarDecl(VD);
return OMPFromClause::Create(
getASTContext(), Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
- MVLI.VarComponents, MVLI.UDMapperList, Modifiers, ModifiersLoc,
- MapperIdScopeSpec.getWithLocInContext(getASTContext()), MapperId);
+ MVLI.VarComponents, MVLI.UDMapperList, IteratorExpr, Modifiers,
+ ModifiersLoc, MapperIdScopeSpec.getWithLocInContext(getASTContext()),
+ MapperId);
}
OMPClause *
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 242ffb09af006..636f54c212f71 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -2196,13 +2196,15 @@ class TreeTransform {
OMPClause *
RebuildOMPToClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
+ Expr *IteratorModifier,
CXXScopeSpec &MapperIdScopeSpec,
DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> UnresolvedMappers) {
return getSema().OpenMP().ActOnOpenMPToClause(
- MotionModifiers, MotionModifiersLoc, MapperIdScopeSpec, MapperId,
- ColonLoc, VarList, Locs, UnresolvedMappers);
+ MotionModifiers, MotionModifiersLoc, IteratorModifier,
+ MapperIdScopeSpec, MapperId, ColonLoc, VarList, Locs,
+ UnresolvedMappers);
}
/// Build a new OpenMP 'from' clause.
@@ -2212,13 +2214,15 @@ class TreeTransform {
OMPClause *
RebuildOMPFromClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
+ Expr *IteratorModifier,
CXXScopeSpec &MapperIdScopeSpec,
DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> UnresolvedMappers) {
return getSema().OpenMP().ActOnOpenMPFromClause(
- MotionModifiers, MotionModifiersLoc, MapperIdScopeSpec, MapperId,
- ColonLoc, VarList, Locs, UnresolvedMappers);
+ MotionModifiers, MotionModifiersLoc, IteratorModifier,
+ MapperIdScopeSpec, MapperId, ColonLoc, VarList, Locs,
+ UnresolvedMappers);
}
/// Build a new OpenMP 'use_device_ptr' clause.
@@ -11411,6 +11415,13 @@ template <typename Derived>
OMPClause *TreeTransform<Derived>::TransformOMPToClause(OMPToClause *C) {
OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
llvm::SmallVector<Expr *, 16> Vars;
+ Expr *IteratorModifier = C->getIteratorModifier();
+ if (IteratorModifier) {
+ ExprResult MapModRes = getDerived().TransformExpr(IteratorModifier);
+ if (MapModRes.isInvalid())
+ return nullptr;
+ IteratorModifier = MapModRes.get();
+ }
CXXScopeSpec MapperIdScopeSpec;
DeclarationNameInfo MapperIdInfo;
llvm::SmallVector<Expr *, 16> UnresolvedMappers;
@@ -11418,14 +11429,22 @@ OMPClause *TreeTransform<Derived>::TransformOMPToClause(OMPToClause *C) {
*this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers))
return nullptr;
return getDerived().RebuildOMPToClause(
- C->getMotionModifiers(), C->getMotionModifiersLoc(), MapperIdScopeSpec,
- MapperIdInfo, C->getColonLoc(), Vars, Locs, UnresolvedMappers);
+ C->getMotionModifiers(), C->getMotionModifiersLoc(), IteratorModifier,
+ MapperIdScopeSpec, MapperIdInfo, C->getColonLoc(), Vars, Locs,
+ UnresolvedMappers);
}
template <typename Derived>
OMPClause *TreeTransform<Derived>::TransformOMPFromClause(OMPFromClause *C) {
OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
llvm::SmallVector<Expr *, 16> Vars;
+ Expr *IteratorModifier = C->getIteratorModifier();
+ if (IteratorModifier) {
+ ExprResult MapModRes = getDerived().TransformExpr(IteratorModifier);
+ if (MapModRes.isInvalid())
+ return nullptr;
+ IteratorModifier = MapModRes.get();
+ }
CXXScopeSpec MapperIdScopeSpec;
DeclarationNameInfo MapperIdInfo;
llvm::SmallVector<Expr *, 16> UnresolvedMappers;
@@ -11433,8 +11452,9 @@ OMPClause *TreeTransform<Derived>::TransformOMPFromClause(OMPFromClause *C) {
*this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers))
return nullptr;
return getDerived().RebuildOMPFromClause(
- C->getMotionModifiers(), C->getMotionModifiersLoc(), MapperIdScopeSpec,
- MapperIdInfo, C->getColonLoc(), Vars, Locs, UnresolvedMappers);
+ C->getMotionModifiers(), C->getMotionModifiersLoc(), IteratorModifier,
+ MapperIdScopeSpec, MapperIdInfo, C->getColonLoc(), Vars, Locs,
+ UnresolvedMappers);
}
template <typename Derived>
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 9ee8a0fb0f060..736ca0e4382bc 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12321,6 +12321,8 @@ void OMPClauseReader::VisitOMPToClause(OMPToClause *C) {
C->setMotionModifier(
I, static_cast<OpenMPMotionModifierKind>(Record.readInt()));
C->setMotionModifierLoc(I, Record.readSourceLocation());
+ if (C->getMotionModifier(I) == OMPC_MOTION_MODIFIER_iterator)
+ C->setIteratorModifier(Record.readExpr());
}
C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc());
C->setMapperIdInfo(Record.readDeclarationNameInfo());
@@ -12377,6 +12379,8 @@ void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
C->setMotionModifier(
I, static_cast<OpenMPMotionModifierKind>(Record.readInt()));
C->setMotionModifierLoc(I, Record.readSourceLocation());
+ if (C->getMotionModifier(I) == OMPC_MOTION_MODIFIER_iterator)
+ C->setIteratorModifier(Record.readExpr());
}
C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc());
C->setMapperIdInfo(Record.readDeclarationNameInfo());
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 2aa77934c08d1..205408ca6e9d2 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8372,6 +8372,8 @@ void OMPClauseWriter::VisitOMPToClause(OMPToClause *C) {
for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
Record.push_back(C->getMotionModifier(I));
Record.AddSourceLocation(C->getMotionModifierLoc(I));
+ if (C->getMotionModifier(I) == OMPC_MOTION_MODIFIER_iterator)
+ Record.AddStmt(C->getIteratorModifier());
}
Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc());
Record.AddDeclarationNameInfo(C->getMapperIdInfo());
@@ -8402,6 +8404,8 @@ void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) {
for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
Record.push_back(C->getMotionModifier(I));
Record.AddSourceLocation(C->getMotionModifierLoc(I));
+ if (C->getMotionModifier(I) == OMPC_MOTION_MODIFIER_iterator)
+ Record.AddStmt(C->getIteratorModifier());
}
Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc());
Record.AddDeclarationNameInfo(C->getMapperIdInfo());
diff --git a/clang/test/OpenMP/target_update_iterator_ast_print.cpp b/clang/test/OpenMP/target_update_iterator_ast_print.cpp
new file mode 100644
index 0000000000000..322f565c9c732
--- /dev/null
+++ b/clang/test/OpenMP/target_update_iterator_ast_print.cpp
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void test() {
+ int a[10];
+ #pragma omp target update to(iterator(int it = 0:10): a[it])
+ // CHECK: int a[10];
+ // CHECK: #pragma omp target update to(iterator(int it = 0:10): a[it])
+ #pragma omp target update from(iterator(int it = 0:10): a[it])
+ // CHECK: #pragma omp target update from(iterator(int it = 0:10): a[it])
+}
+
+#endif
>From d7aa585f082bbe531461f968a73c2cf42eb8584e Mon Sep 17 00:00:00 2001
From: Shashwathi N <nshashwa at pe31.hpc.amslabs.hpecorp.net>
Date: Tue, 16 Sep 2025 10:39:52 -0500
Subject: [PATCH 2/5] Formatting
---
clang/include/clang/AST/OpenMPClause.h | 23 ++++++++++++-----------
clang/lib/AST/OpenMPClause.cpp | 12 ++++++------
clang/lib/Parse/ParseOpenMP.cpp | 24 ++++++++++++------------
clang/lib/Sema/SemaOpenMP.cpp | 18 +++++++++---------
clang/lib/Sema/TreeTransform.h | 8 +++-----
5 files changed, 42 insertions(+), 43 deletions(-)
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 4e5a86482fbd4..9681e35078064 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -7476,7 +7476,8 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
/// Motion-modifiers for the 'to' clause.
OpenMPMotionModifierKind MotionModifiers[NumberOfOMPMotionModifiers] = {
- OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown};
+ OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown,
+ OMPC_MOTION_MODIFIER_unknown};
/// Location of motion-modifiers for the 'to' clause.
SourceLocation MotionModifiersLoc[NumberOfOMPMotionModifiers];
@@ -7585,15 +7586,14 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
/// \param UDMQualifierLoc C++ nested name specifier for the associated
/// user-defined mapper.
/// \param MapperId The identifier of associated user-defined mapper.
- static OMPToClause *Create(const ASTContext &C, const OMPVarListLocTy &Locs,
- ArrayRef<Expr *> Vars,
- ArrayRef<ValueDecl *> Declarations,
- MappableExprComponentListsRef ComponentLists,
- ArrayRef<Expr *> UDMapperRefs,Expr *IteratorModifier,
- ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
- ArrayRef<SourceLocation> MotionModifiersLoc,
- NestedNameSpecifierLoc UDMQualifierLoc,
- DeclarationNameInfo MapperId);
+ static OMPToClause *
+ Create(const ASTContext &C, const OMPVarListLocTy &Locs,
+ ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
+ MappableExprComponentListsRef ComponentLists,
+ ArrayRef<Expr *> UDMapperRefs, Expr *IteratorModifier,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
+ NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId);
/// Creates an empty clause with the place for \a NumVars variables.
///
@@ -7682,7 +7682,8 @@ class OMPFromClause final
/// Motion-modifiers for the 'from' clause.
OpenMPMotionModifierKind MotionModifiers[NumberOfOMPMotionModifiers] = {
- OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown};
+ OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown,
+ OMPC_MOTION_MODIFIER_unknown};
/// Location of motion-modifiers for the 'from' clause.
SourceLocation MotionModifiersLoc[NumberOfOMPMotionModifiers];
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index de352df0f93f9..8146c1fa4d8f4 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -2573,12 +2573,12 @@ template <typename T> void OMPClausePrinter::VisitOMPMotionClause(T *Node) {
if (Node->getMotionModifier(I) == OMPC_MOTION_MODIFIER_iterator)
PrintIterator(OS, Node, Policy);
else {
- OS << getOpenMPSimpleClauseTypeName(Node->getClauseKind(),
- Node->getMotionModifier(I));
- if (Node->getMotionModifier(I) == OMPC_MOTION_MODIFIER_mapper)
- PrintMapper(OS, Node, Policy);
- if (I < ModifierCount - 1)
- OS << ", ";
+ OS << getOpenMPSimpleClauseTypeName(Node->getClauseKind(),
+ Node->getMotionModifier(I));
+ if (Node->getMotionModifier(I) == OMPC_MOTION_MODIFIER_mapper)
+ PrintMapper(OS, Node, Policy);
+ if (I < ModifierCount - 1)
+ OS << ", ";
}
}
}
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 75fd5a12ff3a1..48296c50da5a7 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -4712,19 +4712,19 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
Data.IteratorExpr = Tail.get();
}
} else {
- ConsumeToken();
- if (Modifier == OMPC_MOTION_MODIFIER_mapper) {
- IsInvalidMapperModifier = parseMapperModifier(Data);
- if (IsInvalidMapperModifier)
- break;
- }
- // OpenMP < 5.1 doesn't permit a ',' or additional modifiers.
- if (getLangOpts().OpenMP < 51)
- break;
- // OpenMP 5.1 accepts an optional ',' even if the next character is ':'.
- // TODO: Is that intentional?
- if (Tok.is(tok::comma))
ConsumeToken();
+ if (Modifier == OMPC_MOTION_MODIFIER_mapper) {
+ IsInvalidMapperModifier = parseMapperModifier(Data);
+ if (IsInvalidMapperModifier)
+ break;
+ }
+ // OpenMP < 5.1 doesn't permit a ',' or additional modifiers.
+ if (getLangOpts().OpenMP < 51)
+ break;
+ // OpenMP 5.1 accepts an optional ',' even if the next character is ':'.
+ // TODO: Is that intentional?
+ if (Tok.is(tok::comma))
+ ConsumeToken();
}
}
if (!Data.MotionModifiers.empty() && Tok.isNot(tok::colon)) {
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index f44e6be061c5d..2e37c85a9abab 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -23464,8 +23464,7 @@ void SemaOpenMP::ActOnOpenMPDeclareTargetInitializer(Decl *TargetDecl) {
OMPClause *SemaOpenMP::ActOnOpenMPToClause(
ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
- ArrayRef<SourceLocation> MotionModifiersLoc,
- Expr *IteratorExpr,
+ ArrayRef<SourceLocation> MotionModifiersLoc, Expr *IteratorExpr,
CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
@@ -23494,14 +23493,15 @@ OMPClause *SemaOpenMP::ActOnOpenMPToClause(
MapperIdScopeSpec, MapperId, UnresolvedMappers);
if (MVLI.ProcessedVarList.empty())
return nullptr;
- if(IteratorExpr)
- if (auto *DRE = dyn_cast<DeclRefExpr>(IteratorExpr))
- if (auto *VD = dyn_cast<VarDecl>(DRE->getDecl()))
+ if (IteratorExpr)
+ if (auto *DRE = dyn_cast<DeclRefExpr>(IteratorExpr))
+ if (auto *VD = dyn_cast<VarDecl>(DRE->getDecl()))
DSAStack->addIteratorVarDecl(VD);
return OMPToClause::Create(
getASTContext(), Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
- MVLI.VarComponents, MVLI.UDMapperList,IteratorExpr, Modifiers, ModifiersLoc,
- MapperIdScopeSpec.getWithLocInContext(getASTContext()), MapperId);
+ MVLI.VarComponents, MVLI.UDMapperList, IteratorExpr, Modifiers,
+ ModifiersLoc, MapperIdScopeSpec.getWithLocInContext(getASTContext()),
+ MapperId);
}
OMPClause *SemaOpenMP::ActOnOpenMPFromClause(
@@ -23536,8 +23536,8 @@ OMPClause *SemaOpenMP::ActOnOpenMPFromClause(
if (MVLI.ProcessedVarList.empty())
return nullptr;
if (IteratorExpr)
- if (auto *DRE = dyn_cast<DeclRefExpr>(IteratorExpr))
- if (auto *VD = dyn_cast<VarDecl>(DRE->getDecl()))
+ if (auto *DRE = dyn_cast<DeclRefExpr>(IteratorExpr))
+ if (auto *VD = dyn_cast<VarDecl>(DRE->getDecl()))
DSAStack->addIteratorVarDecl(VD);
return OMPFromClause::Create(
getASTContext(), Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 636f54c212f71..63da489a52164 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -2196,8 +2196,7 @@ class TreeTransform {
OMPClause *
RebuildOMPToClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
- Expr *IteratorModifier,
- CXXScopeSpec &MapperIdScopeSpec,
+ Expr *IteratorModifier, CXXScopeSpec &MapperIdScopeSpec,
DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> UnresolvedMappers) {
@@ -2214,8 +2213,7 @@ class TreeTransform {
OMPClause *
RebuildOMPFromClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
ArrayRef<SourceLocation> MotionModifiersLoc,
- Expr *IteratorModifier,
- CXXScopeSpec &MapperIdScopeSpec,
+ Expr *IteratorModifier, CXXScopeSpec &MapperIdScopeSpec,
DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> UnresolvedMappers) {
@@ -11438,7 +11436,7 @@ template <typename Derived>
OMPClause *TreeTransform<Derived>::TransformOMPFromClause(OMPFromClause *C) {
OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
llvm::SmallVector<Expr *, 16> Vars;
- Expr *IteratorModifier = C->getIteratorModifier();
+ Expr *IteratorModifier = C->getIteratorModifier();
if (IteratorModifier) {
ExprResult MapModRes = getDerived().TransformExpr(IteratorModifier);
if (MapModRes.isInvalid())
>From 98eebb6a44b72b15a19ee2078313469ee895ac2a Mon Sep 17 00:00:00 2001
From: ShashwathiNavada <shashwathinavada at gmail.com>
Date: Wed, 17 Sep 2025 11:17:01 +0530
Subject: [PATCH 3/5] Update ParseOpenMP.cpp
---
clang/lib/Parse/ParseOpenMP.cpp | 4 ++++
1 file changed, 4 insertions(+)
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 48296c50da5a7..790eb23b48b2c 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -4710,6 +4710,10 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
/*DiscardedValue=*/false);
if (Tail.isUsable()) {
Data.IteratorExpr = Tail.get();
+ TPA.Commit();
+ }
+ else
+ TPA.revert();
}
} else {
ConsumeToken();
>From 2ce56c225032db8275627584b08091e7af50e8f0 Mon Sep 17 00:00:00 2001
From: Shashwathi N <nshashwa at pe31.hpc.amslabs.hpecorp.net>
Date: Wed, 17 Sep 2025 01:35:20 -0500
Subject: [PATCH 4/5] Fixed the error
---
clang/lib/Parse/ParseOpenMP.cpp | 6 ++----
1 file changed, 2 insertions(+), 4 deletions(-)
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 790eb23b48b2c..3b1b16ba39ccf 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -4711,10 +4711,8 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
if (Tail.isUsable()) {
Data.IteratorExpr = Tail.get();
TPA.Commit();
- }
- else
- TPA.revert();
- }
+ } else
+ TPA.Revert();
} else {
ConsumeToken();
if (Modifier == OMPC_MOTION_MODIFIER_mapper) {
>From f32b8c81bb7f44cb29ede02b871ecd55f65deee7 Mon Sep 17 00:00:00 2001
From: Shashwathi N <nshashwa at pe31.hpc.amslabs.hpecorp.net>
Date: Tue, 7 Oct 2025 02:47:07 -0500
Subject: [PATCH 5/5] modifications done
---
clang/include/clang/AST/OpenMPClause.h | 4 +-
clang/lib/AST/OpenMPClause.cpp | 4 +-
clang/test/OpenMP/target_update_codegen.cpp | 55 +++++++++++++++++++++
3 files changed, 59 insertions(+), 4 deletions(-)
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 9681e35078064..4fbc44d33caba 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -7614,7 +7614,7 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
"Requested modifier exceeds the total number of modifiers.");
return MotionModifiers[Cnt];
}
- Expr *getIteratorModifier() {
+ Expr *getIteratorModifier() const {
return getTrailingObjects<Expr *>()[2 * varlist_size()];
}
/// Fetches the motion-modifier location at 'Cnt' index of array of modifiers'
@@ -7819,7 +7819,7 @@ class OMPFromClause final
"Requested modifier exceeds the total number of modifiers.");
return MotionModifiers[Cnt];
}
- Expr *getIteratorModifier() {
+ Expr *getIteratorModifier() const {
return getTrailingObjects<Expr *>()[2 * varlist_size()];
}
/// Fetches the motion-modifier location at 'Cnt' index of array of modifiers'
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index 8146c1fa4d8f4..e9127c5052bf1 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -2570,9 +2570,9 @@ template <typename T> void OMPClausePrinter::VisitOMPMotionClause(T *Node) {
OS << '(';
for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
if (Node->getMotionModifier(I) != OMPC_MOTION_MODIFIER_unknown) {
- if (Node->getMotionModifier(I) == OMPC_MOTION_MODIFIER_iterator)
+ if (Node->getMotionModifier(I) == OMPC_MOTION_MODIFIER_iterator) {
PrintIterator(OS, Node, Policy);
- else {
+ } else {
OS << getOpenMPSimpleClauseTypeName(Node->getClauseKind(),
Node->getMotionModifier(I));
if (Node->getMotionModifier(I) == OMPC_MOTION_MODIFIER_mapper)
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
index c8211f475c7fc..2221c78ea197b 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -1560,5 +1560,60 @@ void foo(int arg) {
{ ++arg; }
}
+#endif
+// RUN: %clang_cc1 -DCK26 -verify -Wno-vla -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK26 --check-prefix CK26-64
+// RUN: %clang_cc1 -DCK26 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK26 --check-prefix CK26-64
+// RUN: %clang_cc1 -DCK26 -fopenmp-version=51 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK26 --check-prefix CK26-32
+// RUN: %clang_cc1 -DCK26 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK26 --check-prefix CK26-32
+
+// RUN: %clang_cc1 -DCK26 -verify -Wno-vla -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK26 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK26 -verify -Wno-vla -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK26 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK26
+// CK26: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, ptr }
+// CK26: @[[ITER:[a-zA-Z0-9_]+]] = internal global i32 0, align 4
+void foo(){
+int a[10];
+#pragma omp target update to(iterator(int it = 0:10) : a[it])
+// CHECK: define dso_local void @_Z4testv()
+// CHECK-NEXT: %[[A1:.*]] = alloca \[10 x i32\], align 16
+// CHECK-NEXT: %[[A2:.*]] = alloca \[1 x ptr\], align 8
+// CHECK-NEXT: %[[A3:.*]] = alloca \[1 x ptr\], align 8
+// CHECK-NEXT: %[[A4:.*]] = alloca \[1 x ptr\], align 8
+// CHECK-NEXT: %[[A5:.*]] = alloca \[1 x ptr\], align 8
+// CHECK-NEXT: %[[A6:.*]] = alloca \[1 x ptr\], align 8
+// CHECK-NEXT: %[[A7:.*]] = alloca \[1 x ptr\], align 8
+// CHECK: %[[LOAD1:.*]] = load i32, ptr @_ZZ4testvE16iteratorvariable, align 4
+// CHECK: %[[SEXT1:.*]] = sext i32 %[[LOAD1]] to i64
+// CHECK: %[[GEP1:.*]] = getelementptr inbounds \[10 x i32\], ptr %[[A1]], i64 0, i64 %[[SEXT1]]
+// CHECK: %[[GEP2:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A2]], i32 0, i32 0
+// CHECK: store ptr %[[A1]], ptr %[[GEP2]], align 8
+// CHECK: %[[GEP3:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A3]], i32 0, i32 0
+// CHECK: store ptr %[[GEP1]], ptr %[[GEP3]], align 8
+// CHECK: %[[GEP4:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A4]], i64 0, i64 0
+// CHECK: store ptr null, ptr %[[GEP4]], align 8
+// CHECK: %[[GEP5:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A2]], i32 0, i32 0
+// CHECK: %[[GEP6:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A3]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_update_mapper(ptr @2, i64 -1, i32 1, ptr %[[GEP5]], ptr %[[GEP6]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr @.offload_mapnames, ptr null)
+// CHECK: %[[LOAD2:.*]] = load i32, ptr @[[ITER]], align 4
+// CHECK: %[[SEXT2:.*]] = sext i32 %[[LOAD2]] to i64
+// CHECK: %[[GEP7:.*]] = getelementptr inbounds \[10 x i32\], ptr %[[A1]], i64 0, i64 %[[SEXT2]]
+// CHECK: %[[GEP8:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A5]], i32 0, i32 0
+// CHECK: store ptr %[[A1]], ptr %[[GEP8]], align 8
+// CHECK: %[[GEP9:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A6]], i32 0, i32 0
+// CHECK: store ptr %[[GEP7]], ptr %[[GEP9]], align 8
+// CHECK: %[[GEP10:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A7]], i64 0, i64 0
+// CHECK: store ptr null, ptr %[[GEP10]], align 8
+// CHECK: %[[GEP11:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A5]], i32 0, i32 0
+// CHECK: %[[GEP12:.*]] = getelementptr inbounds \[1 x ptr\], ptr %[[A6]], i32 0, i32 0
+// CHECK: call void @__tgt_target_data_update_mapper(ptr @4, i64 -1, i32 1, ptr %[[GEP11]], ptr %[[GEP12]], ptr @.offload_sizes.1, ptr @.offload_maptypes.2, ptr @.offload_mapnames.3, ptr null)
+// CHECK: ret void
+}
#endif
#endif
More information about the cfe-commits
mailing list