[clang] 9f2f3b9 - [OpenMP] Implement TR8 `present` motion modifier in Clang (1/2)
Joel E. Denny via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 29 09:20:40 PDT 2020
Author: Joel E. Denny
Date: 2020-07-29T12:18:45-04:00
New Revision: 9f2f3b9de6314a009322b6081c792ebf9a469460
URL: https://github.com/llvm/llvm-project/commit/9f2f3b9de6314a009322b6081c792ebf9a469460
DIFF: https://github.com/llvm/llvm-project/commit/9f2f3b9de6314a009322b6081c792ebf9a469460.diff
LOG: [OpenMP] Implement TR8 `present` motion modifier in Clang (1/2)
This patch implements Clang front end support for the OpenMP TR8
`present` motion modifier for `omp target update` directives. The
next patch in this series implements OpenMP runtime support.
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D84711
Added:
Modified:
clang/include/clang/AST/OpenMPClause.h
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Basic/OpenMPKinds.def
clang/include/clang/Basic/OpenMPKinds.h
clang/include/clang/Parse/Parser.h
clang/include/clang/Sema/Sema.h
clang/lib/AST/OpenMPClause.cpp
clang/lib/Basic/OpenMPKinds.cpp
clang/lib/CodeGen/CGOpenMPRuntime.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/OpenMP/declare_mapper_ast_print.c
clang/test/OpenMP/declare_mapper_codegen.cpp
clang/test/OpenMP/target_update_ast_print.cpp
clang/test/OpenMP/target_update_codegen.cpp
clang/test/OpenMP/target_update_messages.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 4f94aa7074ee..5b588f4b5740 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -6329,8 +6329,20 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
friend OMPVarListClause;
friend TrailingObjects;
+ /// Motion-modifiers for the 'to' clause.
+ OpenMPMotionModifierKind MotionModifiers[NumberOfOMPMotionModifiers] = {
+ OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown};
+
+ /// Location of motion-modifiers for the 'to' clause.
+ SourceLocation MotionModifiersLoc[NumberOfOMPMotionModifiers];
+
+ /// Colon location.
+ SourceLocation ColonLoc;
+
/// Build clause with number of variables \a NumVars.
///
+ /// \param TheMotionModifiers Motion-modifiers.
+ /// \param TheMotionModifiersLoc Locations of motion-modifiers.
/// \param MapperQualifierLoc C++ nested name specifier for the associated
/// user-defined mapper.
/// \param MapperIdInfo The identifier of associated user-defined mapper.
@@ -6342,13 +6354,24 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
/// 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 OMPToClause(NestedNameSpecifierLoc MapperQualifierLoc,
+ explicit OMPToClause(ArrayRef<OpenMPMotionModifierKind> TheMotionModifiers,
+ ArrayRef<SourceLocation> TheMotionModifiersLoc,
+ NestedNameSpecifierLoc MapperQualifierLoc,
DeclarationNameInfo MapperIdInfo,
const OMPVarListLocTy &Locs,
const OMPMappableExprListSizeTy &Sizes)
: OMPMappableExprListClause(llvm::omp::OMPC_to, Locs, Sizes,
/*SupportsMapper=*/true, &MapperQualifierLoc,
- &MapperIdInfo) {}
+ &MapperIdInfo) {
+ assert(llvm::array_lengthof(MotionModifiers) == TheMotionModifiers.size() &&
+ "Unexpected number of motion modifiers.");
+ llvm::copy(TheMotionModifiers, std::begin(MotionModifiers));
+
+ assert(llvm::array_lengthof(MotionModifiersLoc) ==
+ TheMotionModifiersLoc.size() &&
+ "Unexpected number of motion modifier locations.");
+ llvm::copy(TheMotionModifiersLoc, std::begin(MotionModifiersLoc));
+ }
/// Build an empty clause.
///
@@ -6361,6 +6384,29 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
: OMPMappableExprListClause(llvm::omp::OMPC_to, OMPVarListLocTy(), Sizes,
/*SupportsMapper=*/true) {}
+ /// Set motion-modifier for the clause.
+ ///
+ /// \param I index for motion-modifier.
+ /// \param T motion-modifier for the clause.
+ void setMotionModifier(unsigned I, OpenMPMotionModifierKind T) {
+ assert(I < NumberOfOMPMotionModifiers &&
+ "Unexpected index to store motion modifier, exceeds array size.");
+ MotionModifiers[I] = T;
+ }
+
+ /// Set location for the motion-modifier.
+ ///
+ /// \param I index for motion-modifier location.
+ /// \param TLoc motion-modifier location.
+ void setMotionModifierLoc(unsigned I, SourceLocation TLoc) {
+ assert(I < NumberOfOMPMotionModifiers &&
+ "Index to store motion modifier location exceeds array size.");
+ MotionModifiersLoc[I] = TLoc;
+ }
+
+ /// Set colon location.
+ void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; }
+
/// 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 {
@@ -6385,6 +6431,8 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
/// \param Vars The original expression used in the clause.
/// \param Declarations Declarations used in the clause.
/// \param ComponentLists Component lists used in the clause.
+ /// \param MotionModifiers Motion-modifiers.
+ /// \param MotionModifiersLoc Location of motion-modifiers.
/// \param UDMapperRefs References to user-defined mappers associated with
/// expressions used in the clause.
/// \param UDMQualifierLoc C++ nested name specifier for the associated
@@ -6395,6 +6443,8 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists,
ArrayRef<Expr *> UDMapperRefs,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
NestedNameSpecifierLoc UDMQualifierLoc,
DeclarationNameInfo MapperId);
@@ -6409,6 +6459,38 @@ class OMPToClause final : public OMPMappableExprListClause<OMPToClause>,
static OMPToClause *CreateEmpty(const ASTContext &C,
const OMPMappableExprListSizeTy &Sizes);
+ /// Fetches the motion-modifier at 'Cnt' index of array of modifiers.
+ ///
+ /// \param Cnt index for motion-modifier.
+ OpenMPMotionModifierKind getMotionModifier(unsigned Cnt) const LLVM_READONLY {
+ assert(Cnt < NumberOfOMPMotionModifiers &&
+ "Requested modifier exceeds the total number of modifiers.");
+ return MotionModifiers[Cnt];
+ }
+
+ /// Fetches the motion-modifier location at 'Cnt' index of array of modifiers'
+ /// locations.
+ ///
+ /// \param Cnt index for motion-modifier location.
+ SourceLocation getMotionModifierLoc(unsigned Cnt) const LLVM_READONLY {
+ assert(Cnt < NumberOfOMPMotionModifiers &&
+ "Requested modifier location exceeds total number of modifiers.");
+ return MotionModifiersLoc[Cnt];
+ }
+
+ /// Fetches ArrayRef of motion-modifiers.
+ ArrayRef<OpenMPMotionModifierKind> getMotionModifiers() const LLVM_READONLY {
+ return llvm::makeArrayRef(MotionModifiers);
+ }
+
+ /// Fetches ArrayRef of location of motion-modifiers.
+ ArrayRef<SourceLocation> getMotionModifiersLoc() const LLVM_READONLY {
+ return llvm::makeArrayRef(MotionModifiersLoc);
+ }
+
+ /// Get colon location.
+ SourceLocation getColonLoc() const { return ColonLoc; }
+
child_range children() {
return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
reinterpret_cast<Stmt **>(varlist_end()));
@@ -6449,8 +6531,20 @@ class OMPFromClause final
friend OMPVarListClause;
friend TrailingObjects;
+ /// Motion-modifiers for the 'from' clause.
+ OpenMPMotionModifierKind MotionModifiers[NumberOfOMPMotionModifiers] = {
+ OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown};
+
+ /// Location of motion-modifiers for the 'from' clause.
+ SourceLocation MotionModifiersLoc[NumberOfOMPMotionModifiers];
+
+ /// Colon location.
+ SourceLocation ColonLoc;
+
/// Build clause with number of variables \a NumVars.
///
+ /// \param TheMotionModifiers Motion-modifiers.
+ /// \param TheMotionModifiersLoc Locations of motion-modifiers.
/// \param MapperQualifierLoc C++ nested name specifier for the associated
/// user-defined mapper.
/// \param MapperIdInfo The identifier of associated user-defined mapper.
@@ -6462,13 +6556,24 @@ class OMPFromClause final
/// 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 OMPFromClause(NestedNameSpecifierLoc MapperQualifierLoc,
+ explicit OMPFromClause(ArrayRef<OpenMPMotionModifierKind> TheMotionModifiers,
+ ArrayRef<SourceLocation> TheMotionModifiersLoc,
+ NestedNameSpecifierLoc MapperQualifierLoc,
DeclarationNameInfo MapperIdInfo,
const OMPVarListLocTy &Locs,
const OMPMappableExprListSizeTy &Sizes)
: OMPMappableExprListClause(llvm::omp::OMPC_from, Locs, Sizes,
/*SupportsMapper=*/true, &MapperQualifierLoc,
- &MapperIdInfo) {}
+ &MapperIdInfo) {
+ assert(llvm::array_lengthof(MotionModifiers) == TheMotionModifiers.size() &&
+ "Unexpected number of motion modifiers.");
+ llvm::copy(TheMotionModifiers, std::begin(MotionModifiers));
+
+ assert(llvm::array_lengthof(MotionModifiersLoc) ==
+ TheMotionModifiersLoc.size() &&
+ "Unexpected number of motion modifier locations.");
+ llvm::copy(TheMotionModifiersLoc, std::begin(MotionModifiersLoc));
+ }
/// Build an empty clause.
///
@@ -6481,6 +6586,29 @@ class OMPFromClause final
: OMPMappableExprListClause(llvm::omp::OMPC_from, OMPVarListLocTy(),
Sizes, /*SupportsMapper=*/true) {}
+ /// Set motion-modifier for the clause.
+ ///
+ /// \param I index for motion-modifier.
+ /// \param T motion-modifier for the clause.
+ void setMotionModifier(unsigned I, OpenMPMotionModifierKind T) {
+ assert(I < NumberOfOMPMotionModifiers &&
+ "Unexpected index to store motion modifier, exceeds array size.");
+ MotionModifiers[I] = T;
+ }
+
+ /// Set location for the motion-modifier.
+ ///
+ /// \param I index for motion-modifier location.
+ /// \param TLoc motion-modifier location.
+ void setMotionModifierLoc(unsigned I, SourceLocation TLoc) {
+ assert(I < NumberOfOMPMotionModifiers &&
+ "Index to store motion modifier location exceeds array size.");
+ MotionModifiersLoc[I] = TLoc;
+ }
+
+ /// Set colon location.
+ void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; }
+
/// 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 {
@@ -6505,18 +6633,21 @@ class OMPFromClause final
/// \param Vars The original expression used in the clause.
/// \param Declarations Declarations used in the clause.
/// \param ComponentLists Component lists used in the clause.
+ /// \param MotionModifiers Motion-modifiers.
+ /// \param MotionModifiersLoc Location of motion-modifiers.
/// \param UDMapperRefs References to user-defined mappers associated with
/// expressions used in the clause.
/// \param UDMQualifierLoc C++ nested name specifier for the associated
/// user-defined mapper.
/// \param MapperId The identifier of associated user-defined mapper.
- static OMPFromClause *Create(const ASTContext &C, const OMPVarListLocTy &Locs,
- ArrayRef<Expr *> Vars,
- ArrayRef<ValueDecl *> Declarations,
- MappableExprComponentListsRef ComponentLists,
- ArrayRef<Expr *> UDMapperRefs,
- NestedNameSpecifierLoc UDMQualifierLoc,
- DeclarationNameInfo MapperId);
+ static OMPFromClause *
+ Create(const ASTContext &C, const OMPVarListLocTy &Locs,
+ ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
+ MappableExprComponentListsRef ComponentLists,
+ ArrayRef<Expr *> UDMapperRefs,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
+ NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId);
/// Creates an empty clause with the place for \a NumVars variables.
///
@@ -6529,6 +6660,38 @@ class OMPFromClause final
static OMPFromClause *CreateEmpty(const ASTContext &C,
const OMPMappableExprListSizeTy &Sizes);
+ /// Fetches the motion-modifier at 'Cnt' index of array of modifiers.
+ ///
+ /// \param Cnt index for motion-modifier.
+ OpenMPMotionModifierKind getMotionModifier(unsigned Cnt) const LLVM_READONLY {
+ assert(Cnt < NumberOfOMPMotionModifiers &&
+ "Requested modifier exceeds the total number of modifiers.");
+ return MotionModifiers[Cnt];
+ }
+
+ /// Fetches the motion-modifier location at 'Cnt' index of array of modifiers'
+ /// locations.
+ ///
+ /// \param Cnt index for motion-modifier location.
+ SourceLocation getMotionModifierLoc(unsigned Cnt) const LLVM_READONLY {
+ assert(Cnt < NumberOfOMPMotionModifiers &&
+ "Requested modifier location exceeds total number of modifiers.");
+ return MotionModifiersLoc[Cnt];
+ }
+
+ /// Fetches ArrayRef of motion-modifiers.
+ ArrayRef<OpenMPMotionModifierKind> getMotionModifiers() const LLVM_READONLY {
+ return llvm::makeArrayRef(MotionModifiers);
+ }
+
+ /// Fetches ArrayRef of location of motion-modifiers.
+ ArrayRef<SourceLocation> getMotionModifiersLoc() const LLVM_READONLY {
+ return llvm::makeArrayRef(MotionModifiersLoc);
+ }
+
+ /// Get colon location.
+ SourceLocation getColonLoc() const { return ColonLoc; }
+
child_range children() {
return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
reinterpret_cast<Stmt **>(varlist_end()));
@@ -7623,6 +7786,8 @@ class OMPClausePrinter final : public OMPClauseVisitor<OMPClausePrinter> {
/// Process clauses with list of variables.
template <typename T> void VisitOMPClauseList(T *Node, char StartSym);
+ /// Process motion clauses.
+ template <typename T> void VisitOMPMotionClause(T *Node);
public:
OMPClausePrinter(raw_ostream &OS, const PrintingPolicy &Policy)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index ae693a08108c..794ef23f0761 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9805,6 +9805,8 @@ def err_omp_multiple_array_items_in_map_clause : Error<
"multiple array elements associated with the same variable are not allowed in map clauses of the same construct">;
def err_omp_duplicate_map_type_modifier : Error<
"same map type modifier has been specified more than once">;
+def err_omp_duplicate_motion_modifier : Error<
+ "same motion modifier has been specified more than once">;
def err_omp_pointer_mapped_along_with_derived_section : Error<
"pointer cannot be mapped along with a section derived from itself">;
def err_omp_original_storage_is_shared_and_does_not_contain : Error<
diff --git a/clang/include/clang/Basic/OpenMPKinds.def b/clang/include/clang/Basic/OpenMPKinds.def
index 04ecbeaaa03e..59bb331a7a8a 100644
--- a/clang/include/clang/Basic/OpenMPKinds.def
+++ b/clang/include/clang/Basic/OpenMPKinds.def
@@ -125,6 +125,7 @@ OPENMP_MAP_MODIFIER_KIND(present)
// Modifiers for 'to' or 'from' clause.
OPENMP_MOTION_MODIFIER_KIND(mapper)
+OPENMP_MOTION_MODIFIER_KIND(present)
// Static attributes for 'dist_schedule' clause.
OPENMP_DIST_SCHEDULE_KIND(static)
diff --git a/clang/include/clang/Basic/OpenMPKinds.h b/clang/include/clang/Basic/OpenMPKinds.h
index 3e9b5c4a8b14..0f37dc9ad997 100644
--- a/clang/include/clang/Basic/OpenMPKinds.h
+++ b/clang/include/clang/Basic/OpenMPKinds.h
@@ -94,6 +94,10 @@ enum OpenMPMotionModifierKind {
OMPC_MOTION_MODIFIER_unknown
};
+/// Number of allowed motion-modifiers.
+static constexpr unsigned NumberOfOMPMotionModifiers =
+ OMPC_MOTION_MODIFIER_unknown;
+
/// OpenMP attributes for 'dist_schedule' clause.
enum OpenMPDistScheduleClauseKind {
#define OPENMP_DIST_SCHEDULE_KIND(Name) OMPC_DIST_SCHEDULE_##Name,
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index e809d87b59a0..4068e6a444c9 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3221,6 +3221,9 @@ class Parser : public CodeCompletionHandler {
MapTypeModifiers;
SmallVector<SourceLocation, NumberOfOMPMapClauseModifiers>
MapTypeModifiersLoc;
+ SmallVector<OpenMPMotionModifierKind, NumberOfOMPMotionModifiers>
+ MotionModifiers;
+ SmallVector<SourceLocation, NumberOfOMPMotionModifiers> MotionModifiersLoc;
bool IsMapTypeImplicit = false;
SourceLocation ExtraModifierLoc;
};
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index fd4300f563a9..4721ccdf1459 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10725,7 +10725,9 @@ class Sema final {
DeclarationNameInfo &ReductionOrMapperId, int ExtraModifier,
ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
ArrayRef<SourceLocation> MapTypeModifiersLoc, bool IsMapTypeImplicit,
- SourceLocation ExtraModifierLoc);
+ SourceLocation ExtraModifierLoc,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc);
/// Called on well-formed 'inclusive' clause.
OMPClause *ActOnOpenMPInclusiveClause(ArrayRef<Expr *> VarList,
SourceLocation StartLoc,
@@ -10862,15 +10864,20 @@ class Sema final {
SourceLocation KindLoc, SourceLocation EndLoc);
/// Called on well-formed 'to' clause.
OMPClause *
- ActOnOpenMPToClause(ArrayRef<Expr *> VarList, CXXScopeSpec &MapperIdScopeSpec,
- DeclarationNameInfo &MapperId,
- const OMPVarListLocTy &Locs,
+ ActOnOpenMPToClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
+ CXXScopeSpec &MapperIdScopeSpec,
+ DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
+ ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
ArrayRef<Expr *> UnresolvedMappers = llvm::None);
/// Called on well-formed 'from' clause.
- OMPClause *ActOnOpenMPFromClause(
- ArrayRef<Expr *> VarList, CXXScopeSpec &MapperIdScopeSpec,
- DeclarationNameInfo &MapperId, const OMPVarListLocTy &Locs,
- ArrayRef<Expr *> UnresolvedMappers = llvm::None);
+ OMPClause *
+ ActOnOpenMPFromClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
+ CXXScopeSpec &MapperIdScopeSpec,
+ DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
+ ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
+ ArrayRef<Expr *> UnresolvedMappers = llvm::None);
/// Called on well-formed 'use_device_ptr' clause.
OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
const OMPVarListLocTy &Locs);
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index 9caa691188fd..e846d325560d 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -1098,6 +1098,8 @@ OMPToClause *OMPToClause::Create(
const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) {
OMPMappableExprListSizeTy Sizes;
Sizes.NumVars = Vars.size();
@@ -1122,7 +1124,8 @@ OMPToClause *OMPToClause::Create(
Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
Sizes.NumComponents));
- auto *Clause = new (Mem) OMPToClause(UDMQualifierLoc, MapperId, Locs, Sizes);
+ auto *Clause = new (Mem) OMPToClause(MotionModifiers, MotionModifiersLoc,
+ UDMQualifierLoc, MapperId, Locs, Sizes);
Clause->setVarRefs(Vars);
Clause->setUDMapperRefs(UDMapperRefs);
@@ -1145,6 +1148,8 @@ OMPFromClause *OMPFromClause::Create(
const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) {
OMPMappableExprListSizeTy Sizes;
Sizes.NumVars = Vars.size();
@@ -1170,7 +1175,8 @@ OMPFromClause *OMPFromClause::Create(
Sizes.NumComponents));
auto *Clause =
- new (Mem) OMPFromClause(UDMQualifierLoc, MapperId, Locs, Sizes);
+ new (Mem) OMPFromClause(MotionModifiers, MotionModifiersLoc,
+ UDMQualifierLoc, MapperId, Locs, Sizes);
Clause->setVarRefs(Vars);
Clause->setUDMapperRefs(UDMapperRefs);
@@ -1937,6 +1943,17 @@ void OMPClausePrinter::VisitOMPDependClause(OMPDependClause *Node) {
OS << ")";
}
+template <typename T>
+static void PrintMapper(raw_ostream &OS, T *Node,
+ const PrintingPolicy &Policy) {
+ OS << '(';
+ NestedNameSpecifier *MapperNNS =
+ Node->getMapperQualifierLoc().getNestedNameSpecifier();
+ if (MapperNNS)
+ MapperNNS->print(OS, Policy);
+ OS << Node->getMapperIdInfo() << ')';
+}
+
void OMPClausePrinter::VisitOMPMapClause(OMPMapClause *Node) {
if (!Node->varlist_empty()) {
OS << "map(";
@@ -1945,14 +1962,8 @@ void OMPClausePrinter::VisitOMPMapClause(OMPMapClause *Node) {
if (Node->getMapTypeModifier(I) != OMPC_MAP_MODIFIER_unknown) {
OS << getOpenMPSimpleClauseTypeName(OMPC_map,
Node->getMapTypeModifier(I));
- if (Node->getMapTypeModifier(I) == OMPC_MAP_MODIFIER_mapper) {
- OS << '(';
- NestedNameSpecifier *MapperNNS =
- Node->getMapperQualifierLoc().getNestedNameSpecifier();
- if (MapperNNS)
- MapperNNS->print(OS, Policy);
- OS << Node->getMapperIdInfo() << ')';
- }
+ if (Node->getMapTypeModifier(I) == OMPC_MAP_MODIFIER_mapper)
+ PrintMapper(OS, Node, Policy);
OS << ',';
}
}
@@ -1964,44 +1975,41 @@ void OMPClausePrinter::VisitOMPMapClause(OMPMapClause *Node) {
}
}
-void OMPClausePrinter::VisitOMPToClause(OMPToClause *Node) {
- if (!Node->varlist_empty()) {
- OS << "to";
- DeclarationNameInfo MapperId = Node->getMapperIdInfo();
- if (MapperId.getName() && !MapperId.getName().isEmpty()) {
- OS << '(';
- OS << "mapper(";
- NestedNameSpecifier *MapperNNS =
- Node->getMapperQualifierLoc().getNestedNameSpecifier();
- if (MapperNNS)
- MapperNNS->print(OS, Policy);
- OS << MapperId << "):";
- VisitOMPClauseList(Node, ' ');
- } else {
- VisitOMPClauseList(Node, '(');
+template <typename T> void OMPClausePrinter::VisitOMPMotionClause(T *Node) {
+ if (Node->varlist_empty())
+ return;
+ OS << getOpenMPClauseName(Node->getClauseKind());
+ unsigned ModifierCount = 0;
+ for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
+ if (Node->getMotionModifier(I) != OMPC_MOTION_MODIFIER_unknown)
+ ++ModifierCount;
+ }
+ if (ModifierCount) {
+ OS << '(';
+ for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
+ if (Node->getMotionModifier(I) != OMPC_MOTION_MODIFIER_unknown) {
+ 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 << ")";
+ OS << ':';
+ VisitOMPClauseList(Node, ' ');
+ } else {
+ VisitOMPClauseList(Node, '(');
}
+ OS << ")";
+}
+
+void OMPClausePrinter::VisitOMPToClause(OMPToClause *Node) {
+ VisitOMPMotionClause(Node);
}
void OMPClausePrinter::VisitOMPFromClause(OMPFromClause *Node) {
- if (!Node->varlist_empty()) {
- OS << "from";
- DeclarationNameInfo MapperId = Node->getMapperIdInfo();
- if (MapperId.getName() && !MapperId.getName().isEmpty()) {
- OS << '(';
- OS << "mapper(";
- NestedNameSpecifier *MapperNNS =
- Node->getMapperQualifierLoc().getNestedNameSpecifier();
- if (MapperNNS)
- MapperNNS->print(OS, Policy);
- OS << MapperId << "):";
- VisitOMPClauseList(Node, ' ');
- } else {
- VisitOMPClauseList(Node, '(');
- }
- OS << ")";
- }
+ VisitOMPMotionClause(Node);
}
void OMPClausePrinter::VisitOMPDistScheduleClause(OMPDistScheduleClause *Node) {
diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp
index da362f99ed29..5c19d60cbd6e 100644
--- a/clang/lib/Basic/OpenMPKinds.cpp
+++ b/clang/lib/Basic/OpenMPKinds.cpp
@@ -64,12 +64,16 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
return Type;
}
case OMPC_to:
- case OMPC_from:
- return llvm::StringSwitch<unsigned>(Str)
+ case OMPC_from: {
+ unsigned Type = llvm::StringSwitch<unsigned>(Str)
#define OPENMP_MOTION_MODIFIER_KIND(Name) \
.Case(#Name, static_cast<unsigned>(OMPC_MOTION_MODIFIER_##Name))
#include "clang/Basic/OpenMPKinds.def"
.Default(OMPC_MOTION_MODIFIER_unknown);
+ if (OpenMPVersion < 51 && Type == OMPC_MOTION_MODIFIER_present)
+ return OMPC_MOTION_MODIFIER_unknown;
+ return Type;
+ }
case OMPC_dist_schedule:
return llvm::StringSwitch<OpenMPDistScheduleClauseKind>(Str)
#define OPENMP_DIST_SCHEDULE_KIND(Name) .Case(#Name, OMPC_DIST_SCHEDULE_##Name)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index f16f2ed89de8..b8e33948c21c 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7122,6 +7122,7 @@ class MappableExprsHandler {
OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
ArrayRef<OpenMPMapModifierKind> MapModifiers;
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers;
bool ReturnDevicePointer = false;
bool IsImplicit = false;
const ValueDecl *Mapper = nullptr;
@@ -7131,10 +7132,12 @@ class MappableExprsHandler {
MapInfo(
OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
OpenMPMapClauseKind MapType,
- ArrayRef<OpenMPMapModifierKind> MapModifiers, bool ReturnDevicePointer,
- bool IsImplicit, const ValueDecl *Mapper = nullptr,
- bool ForDeviceAddr = false)
+ ArrayRef<OpenMPMapModifierKind> MapModifiers,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ bool ReturnDevicePointer, bool IsImplicit,
+ const ValueDecl *Mapper = nullptr, bool ForDeviceAddr = false)
: Components(Components), MapType(MapType), MapModifiers(MapModifiers),
+ MotionModifiers(MotionModifiers),
ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit),
Mapper(Mapper), ForDeviceAddr(ForDeviceAddr) {}
};
@@ -7254,7 +7257,8 @@ class MappableExprsHandler {
/// expression.
OpenMPOffloadMappingFlags getMapTypeBits(
OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
- bool IsImplicit, bool AddPtrFlag, bool AddIsTargetParamFlag) const {
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers, bool IsImplicit,
+ bool AddPtrFlag, bool AddIsTargetParamFlag) const {
OpenMPOffloadMappingFlags Bits =
IsImplicit ? OMP_MAP_IMPLICIT : OMP_MAP_NONE;
switch (MapType) {
@@ -7293,6 +7297,9 @@ class MappableExprsHandler {
if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present)
!= MapModifiers.end())
Bits |= OMP_MAP_PRESENT;
+ if (llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present)
+ != MotionModifiers.end())
+ Bits |= OMP_MAP_PRESENT;
return Bits;
}
@@ -7337,11 +7344,12 @@ class MappableExprsHandler {
/// Generate the base pointers, section pointers, sizes, map type bits, and
/// user-defined mappers (all included in \a CombinedInfo) for the provided
- /// map type, map modifier, and expression components. \a IsFirstComponent
- /// should be set to true if the provided set of components is the first
- /// associated with a capture.
+ /// map type, map or motion modifiers, and expression components.
+ /// \a IsFirstComponent should be set to true if the provided set of
+ /// components is the first associated with a capture.
void generateInfoForComponentList(
OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
MapCombinedInfoTy &CombinedInfo, StructRangeInfoTy &PartialStruct,
bool IsFirstComponentList, bool IsImplicit,
@@ -7694,7 +7702,7 @@ class MappableExprsHandler {
// Emit data for non-overlapped data.
OpenMPOffloadMappingFlags Flags =
OMP_MAP_MEMBER_OF |
- getMapTypeBits(MapType, MapModifiers, IsImplicit,
+ getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit,
/*AddPtrFlag=*/false,
/*AddIsTargetParamFlag=*/false);
LB = BP;
@@ -7750,10 +7758,10 @@ class MappableExprsHandler {
// same expression except for the first one. We also need to signal
// this map is the first one that relates with the current capture
// (there is a set of entries for each capture).
- OpenMPOffloadMappingFlags Flags = getMapTypeBits(
- MapType, MapModifiers, IsImplicit,
- !IsExpressionFirstInfo || RequiresReference,
- IsCaptureFirstInfo && !RequiresReference);
+ OpenMPOffloadMappingFlags Flags =
+ getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit,
+ !IsExpressionFirstInfo || RequiresReference,
+ IsCaptureFirstInfo && !RequiresReference);
if (!IsExpressionFirstInfo) {
// If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
@@ -8017,14 +8025,16 @@ class MappableExprsHandler {
OMPClauseMappableExprCommon::MappableExprComponentListRef L,
OpenMPMapClauseKind MapType,
ArrayRef<OpenMPMapModifierKind> MapModifiers,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
bool ReturnDevicePointer, bool IsImplicit, const ValueDecl *Mapper,
bool ForDeviceAddr = false) {
const ValueDecl *VD =
D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
if (SkipVarSet.count(VD))
return;
- Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
- IsImplicit, Mapper, ForDeviceAddr);
+ Info[VD].emplace_back(L, MapType, MapModifiers, MotionModifiers,
+ ReturnDevicePointer, IsImplicit, Mapper,
+ ForDeviceAddr);
};
assert(CurDir.is<const OMPExecutableDirective *>() &&
@@ -8033,18 +8043,20 @@ class MappableExprsHandler {
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>())
for (const auto L : C->component_lists()) {
InfoGen(std::get<0>(L), std::get<1>(L), C->getMapType(),
- C->getMapTypeModifiers(), /*ReturnDevicePointer=*/false,
- C->isImplicit(), std::get<2>(L));
+ C->getMapTypeModifiers(), llvm::None,
+ /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L));
}
for (const auto *C : CurExecDir->getClausesOfKind<OMPToClause>())
for (const auto L : C->component_lists()) {
InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_to, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L));
+ C->getMotionModifiers(), /*ReturnDevicePointer=*/false,
+ C->isImplicit(), std::get<2>(L));
}
for (const auto *C : CurExecDir->getClausesOfKind<OMPFromClause>())
for (const auto L : C->component_lists()) {
InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_from, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L));
+ C->getMotionModifiers(), /*ReturnDevicePointer=*/false,
+ C->isImplicit(), std::get<2>(L));
}
// Look at the use_device_ptr clause information and mark the existing map
@@ -8097,7 +8109,7 @@ class MappableExprsHandler {
// Nonetheless, generateInfoForComponentList must be called to take
// the pointer into account for the calculation of the range of the
// partial struct.
- InfoGen(nullptr, Components, OMPC_MAP_unknown, llvm::None,
+ InfoGen(nullptr, Components, OMPC_MAP_unknown, llvm::None, llvm::None,
/*ReturnDevicePointer=*/false, C->isImplicit(), nullptr);
DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/false);
} else {
@@ -8162,8 +8174,8 @@ class MappableExprsHandler {
// the pointer into account for the calculation of the range of the
// partial struct.
InfoGen(nullptr, std::get<1>(L), OMPC_MAP_unknown, llvm::None,
- /*ReturnDevicePointer=*/false, C->isImplicit(), nullptr,
- /*ForDeviceAddr=*/true);
+ llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit(),
+ nullptr, /*ForDeviceAddr=*/true);
DeferredInfo[nullptr].emplace_back(IE, VD, /*ForDeviceAddr=*/true);
} else {
llvm::Value *Ptr;
@@ -8195,9 +8207,10 @@ class MappableExprsHandler {
// Remember the current base pointer index.
unsigned CurrentBasePointersIdx = CurInfo.BasePointers.size();
- generateInfoForComponentList(
- L.MapType, L.MapModifiers, L.Components, CurInfo, PartialStruct,
- IsFirstComponentList, L.IsImplicit, L.Mapper, L.ForDeviceAddr);
+ generateInfoForComponentList(L.MapType, L.MapModifiers,
+ L.MotionModifiers, L.Components, CurInfo,
+ PartialStruct, IsFirstComponentList,
+ L.IsImplicit, L.Mapper, L.ForDeviceAddr);
// If this entry relates with a device pointer, set the relevant
// declaration and add the 'return pointer' flag.
@@ -8282,9 +8295,10 @@ class MappableExprsHandler {
std::get<0>(L) ? cast<ValueDecl>(std::get<0>(L)->getCanonicalDecl())
: nullptr;
// Get the corresponding user-defined mapper.
- Info[VD].emplace_back(
- std::get<1>(L), MC->getMapType(), MC->getMapTypeModifiers(),
- /*ReturnDevicePointer=*/false, MC->isImplicit(), std::get<2>(L));
+ Info[VD].emplace_back(std::get<1>(L), MC->getMapType(),
+ MC->getMapTypeModifiers(), llvm::None,
+ /*ReturnDevicePointer=*/false, MC->isImplicit(),
+ std::get<2>(L));
}
}
@@ -8300,9 +8314,10 @@ class MappableExprsHandler {
for (const MapInfo &L : M.second) {
assert(!L.Components.empty() &&
"Not expecting declaration with no component lists.");
- generateInfoForComponentList(
- L.MapType, L.MapModifiers, L.Components, CurInfo, PartialStruct,
- IsFirstComponentList, L.IsImplicit, L.Mapper, L.ForDeviceAddr);
+ generateInfoForComponentList(L.MapType, L.MapModifiers,
+ L.MotionModifiers, L.Components, CurInfo,
+ PartialStruct, IsFirstComponentList,
+ L.IsImplicit, L.Mapper, L.ForDeviceAddr);
IsFirstComponentList = false;
}
@@ -8568,9 +8583,9 @@ class MappableExprsHandler {
OverlappedComponents = Pair.getSecond();
bool IsFirstComponentList = true;
generateInfoForComponentList(
- MapType, MapModifiers, Components, CombinedInfo, PartialStruct,
- IsFirstComponentList, IsImplicit, Mapper, /*ForDeviceAddr=*/false,
- OverlappedComponents);
+ MapType, MapModifiers, llvm::None, Components, CombinedInfo,
+ PartialStruct, IsFirstComponentList, IsImplicit, Mapper,
+ /*ForDeviceAddr=*/false, OverlappedComponents);
}
// Go through other elements without overlapped elements.
bool IsFirstComponentList = OverlappedData.empty();
@@ -8583,8 +8598,8 @@ class MappableExprsHandler {
std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L;
auto It = OverlappedData.find(&L);
if (It == OverlappedData.end())
- generateInfoForComponentList(MapType, MapModifiers, Components,
- CombinedInfo, PartialStruct,
+ generateInfoForComponentList(MapType, MapModifiers, llvm::None,
+ Components, CombinedInfo, PartialStruct,
IsFirstComponentList, IsImplicit, Mapper);
IsFirstComponentList = false;
}
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index ac1366e6b0ef..8a92a8259c44 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -3453,25 +3453,43 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
if (Tok.is(tok::colon))
Data.ColonLoc = ConsumeToken();
} else if (Kind == OMPC_to || Kind == OMPC_from) {
- if (Tok.is(tok::identifier)) {
+ while (Tok.is(tok::identifier)) {
auto Modifier =
static_cast<OpenMPMotionModifierKind>(getOpenMPSimpleClauseType(
Kind, PP.getSpelling(Tok), getLangOpts().OpenMP));
+ if (Modifier == OMPC_MOTION_MODIFIER_unknown)
+ break;
+ Data.MotionModifiers.push_back(Modifier);
+ Data.MotionModifiersLoc.push_back(Tok.getLocation());
+ ConsumeToken();
if (Modifier == OMPC_MOTION_MODIFIER_mapper) {
- // Parse the mapper modifier.
- ConsumeToken();
IsInvalidMapperModifier = parseMapperModifier(Data);
- if (Tok.isNot(tok::colon)) {
- if (!IsInvalidMapperModifier)
- Diag(Tok, diag::warn_pragma_expected_colon) << ")";
- SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end,
- StopBeforeMatch);
- }
- // Consume ':'.
- if (Tok.is(tok::colon))
- ConsumeToken();
+ 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)) {
+ if (!IsInvalidMapperModifier) {
+ if (getLangOpts().OpenMP < 51)
+ Diag(Tok, diag::warn_pragma_expected_colon) << ")";
+ else
+ Diag(Tok, diag::warn_pragma_expected_colon) << "motion modifier";
+ }
+ SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end,
+ StopBeforeMatch);
}
+ // OpenMP 5.1 permits a ':' even without a preceding modifier. TODO: Is
+ // that intentional?
+ if ((!Data.MotionModifiers.empty() || getLangOpts().OpenMP >= 51) &&
+ Tok.is(tok::colon))
+ Data.ColonLoc = ConsumeToken();
} else if (Kind == OMPC_allocate ||
(Kind == OMPC_affinity && Tok.is(tok::identifier) &&
PP.getSpelling(Tok) == "iterator")) {
@@ -3647,6 +3665,7 @@ OMPClause *Parser::ParseOpenMPVarListClause(OpenMPDirectiveKind DKind,
Kind, Vars, Data.DepModOrTailExpr, Locs, Data.ColonLoc,
Data.ReductionOrMapperIdScopeSpec, Data.ReductionOrMapperId,
Data.ExtraModifier, Data.MapTypeModifiers, Data.MapTypeModifiersLoc,
- Data.IsMapTypeImplicit, Data.ExtraModifierLoc);
+ Data.IsMapTypeImplicit, Data.ExtraModifierLoc, Data.MotionModifiers,
+ Data.MotionModifiersLoc);
}
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 4f69975c54ed..0367d5c8c601 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -13544,7 +13544,9 @@ OMPClause *Sema::ActOnOpenMPVarListClause(
DeclarationNameInfo &ReductionOrMapperId, int ExtraModifier,
ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
ArrayRef<SourceLocation> MapTypeModifiersLoc, bool IsMapTypeImplicit,
- SourceLocation ExtraModifierLoc) {
+ SourceLocation ExtraModifierLoc,
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc) {
SourceLocation StartLoc = Locs.StartLoc;
SourceLocation LParenLoc = Locs.LParenLoc;
SourceLocation EndLoc = Locs.EndLoc;
@@ -13621,12 +13623,14 @@ OMPClause *Sema::ActOnOpenMPVarListClause(
IsMapTypeImplicit, ExtraModifierLoc, ColonLoc, VarList, Locs);
break;
case OMPC_to:
- Res = ActOnOpenMPToClause(VarList, ReductionOrMapperIdScopeSpec,
- ReductionOrMapperId, Locs);
+ Res = ActOnOpenMPToClause(MotionModifiers, MotionModifiersLoc,
+ ReductionOrMapperIdScopeSpec, ReductionOrMapperId,
+ ColonLoc, VarList, Locs);
break;
case OMPC_from:
- Res = ActOnOpenMPFromClause(VarList, ReductionOrMapperIdScopeSpec,
- ReductionOrMapperId, Locs);
+ Res = ActOnOpenMPFromClause(MotionModifiers, MotionModifiersLoc,
+ ReductionOrMapperIdScopeSpec,
+ ReductionOrMapperId, ColonLoc, VarList, Locs);
break;
case OMPC_use_device_ptr:
Res = ActOnOpenMPUseDevicePtrClause(VarList, Locs);
@@ -18584,11 +18588,31 @@ void Sema::checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D,
checkDeclInTargetContext(E->getExprLoc(), E->getSourceRange(), *this, D);
}
-OMPClause *Sema::ActOnOpenMPToClause(ArrayRef<Expr *> VarList,
- CXXScopeSpec &MapperIdScopeSpec,
- DeclarationNameInfo &MapperId,
- const OMPVarListLocTy &Locs,
- ArrayRef<Expr *> UnresolvedMappers) {
+OMPClause *Sema::ActOnOpenMPToClause(
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
+ 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};
+ SourceLocation ModifiersLoc[NumberOfOMPMotionModifiers];
+
+ // Process motion-modifiers, flag errors for duplicate modifiers.
+ unsigned Count = 0;
+ for (unsigned I = 0, E = MotionModifiers.size(); I < E; ++I) {
+ if (MotionModifiers[I] != OMPC_MOTION_MODIFIER_unknown &&
+ llvm::find(Modifiers, MotionModifiers[I]) != std::end(Modifiers)) {
+ Diag(MotionModifiersLoc[I], diag::err_omp_duplicate_motion_modifier);
+ continue;
+ }
+ assert(Count < NumberOfOMPMotionModifiers &&
+ "Modifiers exceed the allowed number of motion modifiers");
+ Modifiers[Count] = MotionModifiers[I];
+ ModifiersLoc[Count] = MotionModifiersLoc[I];
+ ++Count;
+ }
+
MappableVarListInfo MVLI(VarList);
checkMappableExpressionList(*this, DSAStack, OMPC_to, MVLI, Locs.StartLoc,
MapperIdScopeSpec, MapperId, UnresolvedMappers);
@@ -18597,15 +18621,35 @@ OMPClause *Sema::ActOnOpenMPToClause(ArrayRef<Expr *> VarList,
return OMPToClause::Create(
Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
- MVLI.VarComponents, MVLI.UDMapperList,
+ MVLI.VarComponents, MVLI.UDMapperList, Modifiers, ModifiersLoc,
MapperIdScopeSpec.getWithLocInContext(Context), MapperId);
}
-OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef<Expr *> VarList,
- CXXScopeSpec &MapperIdScopeSpec,
- DeclarationNameInfo &MapperId,
- const OMPVarListLocTy &Locs,
- ArrayRef<Expr *> UnresolvedMappers) {
+OMPClause *Sema::ActOnOpenMPFromClause(
+ ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
+ 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};
+ SourceLocation ModifiersLoc[NumberOfOMPMotionModifiers];
+
+ // Process motion-modifiers, flag errors for duplicate modifiers.
+ unsigned Count = 0;
+ for (unsigned I = 0, E = MotionModifiers.size(); I < E; ++I) {
+ if (MotionModifiers[I] != OMPC_MOTION_MODIFIER_unknown &&
+ llvm::find(Modifiers, MotionModifiers[I]) != std::end(Modifiers)) {
+ Diag(MotionModifiersLoc[I], diag::err_omp_duplicate_motion_modifier);
+ continue;
+ }
+ assert(Count < NumberOfOMPMotionModifiers &&
+ "Modifiers exceed the allowed number of motion modifiers");
+ Modifiers[Count] = MotionModifiers[I];
+ ModifiersLoc[Count] = MotionModifiersLoc[I];
+ ++Count;
+ }
+
MappableVarListInfo MVLI(VarList);
checkMappableExpressionList(*this, DSAStack, OMPC_from, MVLI, Locs.StartLoc,
MapperIdScopeSpec, MapperId, UnresolvedMappers);
@@ -18614,7 +18658,7 @@ OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef<Expr *> VarList,
return OMPFromClause::Create(
Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
- MVLI.VarComponents, MVLI.UDMapperList,
+ MVLI.VarComponents, MVLI.UDMapperList, Modifiers, ModifiersLoc,
MapperIdScopeSpec.getWithLocInContext(Context), MapperId);
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 0f83431c94cf..53064d28dd99 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -2006,26 +2006,32 @@ class TreeTransform {
///
/// By default, performs semantic analysis to build the new statement.
/// Subclasses may override this routine to provide
diff erent behavior.
- OMPClause *RebuildOMPToClause(ArrayRef<Expr *> VarList,
- CXXScopeSpec &MapperIdScopeSpec,
- DeclarationNameInfo &MapperId,
- const OMPVarListLocTy &Locs,
- ArrayRef<Expr *> UnresolvedMappers) {
- return getSema().ActOnOpenMPToClause(VarList, MapperIdScopeSpec, MapperId,
- Locs, UnresolvedMappers);
+ OMPClause *
+ RebuildOMPToClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
+ CXXScopeSpec &MapperIdScopeSpec,
+ DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
+ ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
+ ArrayRef<Expr *> UnresolvedMappers) {
+ return getSema().ActOnOpenMPToClause(MotionModifiers, MotionModifiersLoc,
+ MapperIdScopeSpec, MapperId, ColonLoc,
+ VarList, Locs, UnresolvedMappers);
}
/// Build a new OpenMP 'from' clause.
///
/// By default, performs semantic analysis to build the new statement.
/// Subclasses may override this routine to provide
diff erent behavior.
- OMPClause *RebuildOMPFromClause(ArrayRef<Expr *> VarList,
- CXXScopeSpec &MapperIdScopeSpec,
- DeclarationNameInfo &MapperId,
- const OMPVarListLocTy &Locs,
- ArrayRef<Expr *> UnresolvedMappers) {
- return getSema().ActOnOpenMPFromClause(VarList, MapperIdScopeSpec, MapperId,
- Locs, UnresolvedMappers);
+ OMPClause *
+ RebuildOMPFromClause(ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
+ ArrayRef<SourceLocation> MotionModifiersLoc,
+ CXXScopeSpec &MapperIdScopeSpec,
+ DeclarationNameInfo &MapperId, SourceLocation ColonLoc,
+ ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
+ ArrayRef<Expr *> UnresolvedMappers) {
+ return getSema().ActOnOpenMPFromClause(
+ MotionModifiers, MotionModifiersLoc, MapperIdScopeSpec, MapperId,
+ ColonLoc, VarList, Locs, UnresolvedMappers);
}
/// Build a new OpenMP 'use_device_ptr' clause.
@@ -9738,8 +9744,9 @@ OMPClause *TreeTransform<Derived>::TransformOMPToClause(OMPToClause *C) {
if (transformOMPMappableExprListClause<Derived, OMPToClause>(
*this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers))
return nullptr;
- return getDerived().RebuildOMPToClause(Vars, MapperIdScopeSpec, MapperIdInfo,
- Locs, UnresolvedMappers);
+ return getDerived().RebuildOMPToClause(
+ C->getMotionModifiers(), C->getMotionModifiersLoc(), MapperIdScopeSpec,
+ MapperIdInfo, C->getColonLoc(), Vars, Locs, UnresolvedMappers);
}
template <typename Derived>
@@ -9753,7 +9760,8 @@ OMPClause *TreeTransform<Derived>::TransformOMPFromClause(OMPFromClause *C) {
*this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers))
return nullptr;
return getDerived().RebuildOMPFromClause(
- Vars, MapperIdScopeSpec, MapperIdInfo, Locs, UnresolvedMappers);
+ C->getMotionModifiers(), C->getMotionModifiersLoc(), 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 6b96b4ff59b8..a27c36420fdb 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12601,8 +12601,14 @@ void OMPClauseReader::VisitOMPDefaultmapClause(OMPDefaultmapClause *C) {
void OMPClauseReader::VisitOMPToClause(OMPToClause *C) {
C->setLParenLoc(Record.readSourceLocation());
+ for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
+ C->setMotionModifier(
+ I, static_cast<OpenMPMotionModifierKind>(Record.readInt()));
+ C->setMotionModifierLoc(I, Record.readSourceLocation());
+ }
C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc());
C->setMapperIdInfo(Record.readDeclarationNameInfo());
+ C->setColonLoc(Record.readSourceLocation());
auto NumVars = C->varlist_size();
auto UniqueDecls = C->getUniqueDeclarationsNum();
auto TotalLists = C->getTotalComponentListNum();
@@ -12651,8 +12657,14 @@ void OMPClauseReader::VisitOMPToClause(OMPToClause *C) {
void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
C->setLParenLoc(Record.readSourceLocation());
+ for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
+ C->setMotionModifier(
+ I, static_cast<OpenMPMotionModifierKind>(Record.readInt()));
+ C->setMotionModifierLoc(I, Record.readSourceLocation());
+ }
C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc());
C->setMapperIdInfo(Record.readDeclarationNameInfo());
+ C->setColonLoc(Record.readSourceLocation());
auto NumVars = C->varlist_size();
auto UniqueDecls = C->getUniqueDeclarationsNum();
auto TotalLists = C->getTotalComponentListNum();
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 9ea90b2a0212..8c693b7d34fd 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -6581,8 +6581,13 @@ void OMPClauseWriter::VisitOMPToClause(OMPToClause *C) {
Record.push_back(C->getTotalComponentListNum());
Record.push_back(C->getTotalComponentsNum());
Record.AddSourceLocation(C->getLParenLoc());
+ for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
+ Record.push_back(C->getMotionModifier(I));
+ Record.AddSourceLocation(C->getMotionModifierLoc(I));
+ }
Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc());
Record.AddDeclarationNameInfo(C->getMapperIdInfo());
+ Record.AddSourceLocation(C->getColonLoc());
for (auto *E : C->varlists())
Record.AddStmt(E);
for (auto *E : C->mapperlists())
@@ -6605,8 +6610,13 @@ void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) {
Record.push_back(C->getTotalComponentListNum());
Record.push_back(C->getTotalComponentsNum());
Record.AddSourceLocation(C->getLParenLoc());
+ for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
+ Record.push_back(C->getMotionModifier(I));
+ Record.AddSourceLocation(C->getMotionModifierLoc(I));
+ }
Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc());
Record.AddDeclarationNameInfo(C->getMapperIdInfo());
+ Record.AddSourceLocation(C->getColonLoc());
for (auto *E : C->varlists())
Record.AddStmt(E);
for (auto *E : C->mapperlists())
diff --git a/clang/test/OpenMP/declare_mapper_ast_print.c b/clang/test/OpenMP/declare_mapper_ast_print.c
index e82bc4c2d310..b7c9976c122a 100644
--- a/clang/test/OpenMP/declare_mapper_ast_print.c
+++ b/clang/test/OpenMP/declare_mapper_ast_print.c
@@ -5,6 +5,14 @@
// RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s
// RUN: %clang_cc1 -fopenmp-simd -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
// expected-no-diagnostics
#ifndef HEADER
@@ -50,6 +58,12 @@ int main() {
{ dd[0].i++; }
#pragma omp target update to(mapper(id): vv) from(mapper(default): dd[0:10])
// CHECK: #pragma omp target update to(mapper(id): vv) from(mapper(default): dd[0:10])
+#ifdef OMP51
+#pragma omp target update to(mapper(id) present: vv) from(mapper(default), present: dd[0:10])
+// OMP51: #pragma omp target update to(mapper(id), present: vv) from(mapper(default), present: dd[0:10])
+#pragma omp target update to(present mapper(id): vv) from(present, mapper(default): dd[0:10])
+// OMP51: #pragma omp target update to(present, mapper(id): vv) from(present, mapper(default): dd[0:10])
+#endif
}
return 0;
}
diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp
index ee64fe209900..2fd4b3cb7ed0 100644
--- a/clang/test/OpenMP/declare_mapper_codegen.cpp
+++ b/clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -839,4 +839,242 @@ void foo(int a){
#endif // CK3
+///==========================================================================///
+// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK4 --check-prefix CK4-64 %s
+// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK4 --check-prefix CK4-64 %s
+// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK4 --check-prefix CK4-32 %s
+// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK4 --check-prefix CK4-32 %s
+
+// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
+// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+
+#ifdef CK4
+// Mapper function code generation and runtime interface.
+
+// CK4-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
+// CK4-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
+// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
+// CK4: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+
+// CK4-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
+// CK4-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
+// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 = 0x1022
+// CK4: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1022]]]
+
+class C {
+public:
+ int a;
+ double *b;
+};
+
+#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
+
+// CK4: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}})
+// CK4: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]]
+// CK4: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]]
+// CK4: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]]
+// CK4: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]]
+// CK4: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]]
+// CK4-DAG: [[BYTESIZE:%.+]] = load i64, i64* [[SIZEADDR]]
+// CK4-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
+// CK4-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8
+// CK4-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]]
+// CK4-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]]
+// CK4-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C**
+// CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]]
+// CK4-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]]
+// CK4-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]]
+// CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
+// CK4: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
+
+// CK4: [[INITEVALDEL]]
+// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK4: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
+// CK4: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
+// CK4: [[INIT]]
+// CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
+// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
+// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]])
+// CK4: br label %[[LHEAD:[^,]+]]
+
+// CK4: [[LHEAD]]
+// CK4: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]]
+// CK4: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
+// CK4: [[LBODY]]
+// CK4: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
+// CK4: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]]
+// CK4-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0
+// CK4-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
+// CK4-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1
+// CK4-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]]
+// CK4-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz:64|32]] 0
+// CK4-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1
+// CK4-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8*
+// CK4-DAG: [[BENDV:%.+]] = bitcast double** [[BEND]] to i8*
+// CK4-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64
+// CK4-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64
+// CK4-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
+// CK4-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+// CK4-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
+// CK4-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
+// CK4-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]])
+// CK4-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
+// CK4-DAG: br label %[[MEMBER:[^,]+]]
+// CK4-DAG: [[MEMBER]]
+// CK4-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
+// CK4-DAG: [[MEMBERCOM]]
+// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK4-DAG: br label %[[LTYPE]]
+// CK4-DAG: [[LTYPE]]
+// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK4-DAG: [[ALLOC]]
+// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK4-DAG: br label %[[TYEND:[^,]+]]
+// CK4-DAG: [[ALLOCELSE]]
+// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK4-DAG: [[TO]]
+// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
+// CK4-DAG: br label %[[TYEND]]
+// CK4-DAG: [[TOELSE]]
+// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK4-DAG: [[FROM]]
+// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
+// CK4-DAG: br label %[[TYEND]]
+// CK4-DAG: [[TYEND]]
+// CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]])
+// CK4-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8*
+// CK4-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8*
+// CK4-DAG: br label %[[MEMBER:[^,]+]]
+// CK4-DAG: [[MEMBER]]
+// CK4-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
+// CK4-DAG: [[MEMBERCOM]]
+// 281474976710659 == 0x1,000,000,003
+// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
+// CK4-DAG: br label %[[LTYPE]]
+// CK4-DAG: [[LTYPE]]
+// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK4-DAG: [[ALLOC]]
+// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK4-DAG: br label %[[TYEND:[^,]+]]
+// CK4-DAG: [[ALLOCELSE]]
+// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK4-DAG: [[TO]]
+// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
+// CK4-DAG: br label %[[TYEND]]
+// CK4-DAG: [[TOELSE]]
+// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK4-DAG: [[FROM]]
+// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
+// CK4-DAG: br label %[[TYEND]]
+// CK4-DAG: [[TYEND]]
+// CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]])
+// CK4-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8*
+// CK4-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8*
+// CK4-DAG: br label %[[MEMBER:[^,]+]]
+// CK4-DAG: [[MEMBER]]
+// CK4-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
+// CK4-DAG: [[MEMBERCOM]]
+// 281474976710675 == 0x1,000,000,013
+// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
+// CK4-DAG: br label %[[LTYPE]]
+// CK4-DAG: [[LTYPE]]
+// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710675, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
+// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
+// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
+// CK4-DAG: [[ALLOC]]
+// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
+// CK4-DAG: br label %[[TYEND:[^,]+]]
+// CK4-DAG: [[ALLOCELSE]]
+// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
+// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
+// CK4-DAG: [[TO]]
+// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
+// CK4-DAG: br label %[[TYEND]]
+// CK4-DAG: [[TOELSE]]
+// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
+// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
+// CK4-DAG: [[FROM]]
+// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
+// CK4-DAG: br label %[[TYEND]]
+// CK4-DAG: [[TYEND]]
+// CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]])
+// CK4: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1
+// CK4: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]]
+// CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
+
+// CK4: [[LEXIT]]
+// CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1
+// CK4: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
+// CK4: [[EVALDEL]]
+// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
+// CK4: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
+// CK4: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]]
+// CK4: [[DEL]]
+// CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
+// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
+// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
+// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]])
+// CK4: br label %[[DONE]]
+// CK4: [[DONE]]
+// CK4: ret void
+
+
+// CK4-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
+void foo(int a){
+ int i = a;
+ C c;
+ c.a = a;
+
+ // CK4-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}, i8** [[TMPRGEP:%.+]])
+ // CK4-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK4-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK4-DAG: [[TMPRGEP]] = bitcast [1 x i8*]* [[TMPR:%[^,]+]] to i8**
+ // CK4-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
+ // CK4-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
+ // CK4-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i[[sz]] 0, i[[sz]] 0
+ // CK4-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C**
+ // CK4-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
+ // CK4-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[TCBP0]]
+ // CK4-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
+ // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[TMPR1]]
+ #pragma omp target update to(present, mapper(id): c)
+
+ // CK4-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}, i8** [[FMPRGEP:%.+]])
+ // CK4-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK4-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
+ // CK4-DAG: [[FMPRGEP]] = bitcast [1 x i8*]* [[FMPR:%[^,]+]] to i8**
+ // CK4-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
+ // CK4-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
+ // CK4-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 0
+ // CK4-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
+ // CK4-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
+ // CK4-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
+ // CK4-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
+ // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
+ #pragma omp target update from(mapper(id), present: c)
+}
+
+#endif // CK4
+
#endif // HEADER
diff --git a/clang/test/OpenMP/target_update_ast_print.cpp b/clang/test/OpenMP/target_update_ast_print.cpp
index 0111432fde8f..745233c430eb 100644
--- a/clang/test/OpenMP/target_update_ast_print.cpp
+++ b/clang/test/OpenMP/target_update_ast_print.cpp
@@ -5,6 +5,14 @@
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
// expected-no-diagnostics
#ifndef HEADER
@@ -25,6 +33,13 @@ T foo(T targ, U uarg) {
#pragma omp target update to(arr[2][0:1:2])
#pragma omp target update from(arr[2][0:1:2])
+
+#ifdef OMP51
+#pragma omp target update to(present: arr[2][0:1:2])
+
+#pragma omp target update from(present: arr[2][0:1:2], a)
+#endif
+
return a + targ + (T)b;
}
// CHECK: static T a, *p;
@@ -45,6 +60,8 @@ T foo(T targ, U uarg) {
// CHECK: int arr[100][100];
// CHECK-NEXT: #pragma omp target update to(arr[2][0:1:2])
// CHECK-NEXT: #pragma omp target update from(arr[2][0:1:2])
+// OMP5-NEXT: #pragma omp target update to(present: arr[2][0:1:2])
+// OMP5-NEXT: #pragma omp target update from(present: arr[2][0:1:2], a)
int main(int argc, char **argv) {
static int a;
@@ -62,6 +79,13 @@ int main(int argc, char **argv) {
// CHECK-NEXT: #pragma omp target update to(argv[2][0:1:2])
#pragma omp target update from(argv[2][0:1:2])
// CHECK-NEXT: #pragma omp target update from(argv[2][0:1:2])
+#ifdef OMP51
+#pragma omp target update to(present: argv[2][0:1:2])
+// OMP5-NEXT: #pragma omp target update to(present: arr[2][0:1:2])
+#pragma omp target update from(argv[2][0:1:2], a)
+// OMP5-NEXT: #pragma omp target update from(present: arr[2][0:1:2], a)
+#endif
+
return foo(argc, f) + foo(argv[0][0], f) + a;
}
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
index 637b580d754f..feea1129be4c 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -1084,5 +1084,68 @@ void array_shaping(float *f, int sa) {
#pragma omp target update from(([sa][5])f)
}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64
+// RUN: %clang_cc1 -DCK19 -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 %s -emit-llvm -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-32
+// RUN: %clang_cc1 -DCK19 -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 %s -emit-llvm -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-32
+
+// RUN: %clang_cc1 -DCK19 -verify -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-ONLY0 %s
+// RUN: %clang_cc1 -DCK19 -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 %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK19 -verify -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK19 -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 %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK19
+
+// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
+// CK19: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+
+// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 = 0x1022
+// CK19: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1022]]]
+
+// CK19-LABEL: _Z13check_presenti
+void check_present(int arg) {
+ int la;
+ float lb[arg];
+
+ // Region 00
+ // CK19-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+ // CK19-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
+ // CK19-DAG: store float* [[VAL0]], float** [[CP0]]
+ // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
+ #pragma omp target update to(present: lb)
+ ;
+
+ // Region 01
+ // CK19-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}, i8** null)
+ // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+ // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+ // CK19-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
+ // CK19-DAG: store float* [[VAL0]], float** [[CP0]]
+ // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
+ #pragma omp target update from(present: lb)
+ ;
+}
#endif
#endif
diff --git a/clang/test/OpenMP/target_update_messages.cpp b/clang/test/OpenMP/target_update_messages.cpp
index 4092b623bdc9..401941b63028 100644
--- a/clang/test/OpenMP/target_update_messages.cpp
+++ b/clang/test/OpenMP/target_update_messages.cpp
@@ -1,8 +1,10 @@
-// RUN: %clang_cc1 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp50 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
void xxx(int argc) {
int x; // expected-note {{initialize the variable 'x' to silence this warning}}
@@ -25,6 +27,10 @@ int tmain(T argc, S **argv) {
return 0;
}
+struct S {
+ int i;
+};
+
int main(int argc, char **argv) {
int m;
#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
@@ -34,15 +40,114 @@ int main(int argc, char **argv) {
#pragma omp target update to(m) ] // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
#pragma omp target update to(m) ) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+ #pragma omp declare mapper(id: S s) map(s.i)
+ S s;
+
+ // Check parsing with no modifiers.
+ // lt51-error at +2 {{expected expression}}
+ // lt51-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(: s)
+ // expected-error at +2 {{expected expression}}
+ // expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(:)
+ // expected-error at +2 2 {{expected expression}}
+ // expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(,:)
+
+ // Check parsing with one modifier.
+ // expected-error at +2 {{use of undeclared identifier 'foobar'}}
+ // expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(foobar: s)
+ // expected-error at +3 {{expected ',' or ')' in 'to' clause}}
+ // expected-error at +2 {{expected ')'}}
+ // expected-note at +1 {{to match this '('}}
+ #pragma omp target update to(m: s)
+ #pragma omp target update to(mapper(id): s)
+ // lt51-error at +2 {{use of undeclared identifier 'present'}}
+ // lt51-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(present: s)
+ // ge51-warning at +4 {{missing ':' after motion modifier - ignoring}}
+ // lt51-warning at +3 {{missing ':' after ) - ignoring}}
+ // expected-error at +2 {{expected expression}}
+ // expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(mapper(id) s)
+ // ge51-warning at +4 {{missing ':' after motion modifier - ignoring}}
+ // ge51-error at +3 {{expected expression}}
+ // lt51-error at +2 {{use of undeclared identifier 'present'}}
+ // expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(present s)
+ // ge51-warning at +4 {{missing ':' after motion modifier - ignoring}}
+ // lt51-warning at +3 {{missing ':' after ) - ignoring}}
+ // expected-error at +2 {{expected expression}}
+ // expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(mapper(id))
+ // ge51-warning at +4 {{missing ':' after motion modifier - ignoring}}
+ // ge51-error at +3 {{expected expression}}
+ // lt51-error at +2 {{use of undeclared identifier 'present'}}
+ // expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(present)
+ // expected-error at +2 {{expected expression}}
+ // expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(mapper(id):)
+ // ge51-error at +3 {{expected expression}}
+ // lt51-error at +2 {{use of undeclared identifier 'present'}}
+ // expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(present:)
+
+ // Check parsing with two modifiers.
+ // lt51-warning at +1 {{missing ':' after ) - ignoring}}
+ #pragma omp target update to(mapper(id), present: s)
+ // lt51-error at +3 {{use of undeclared identifier 'present'}}
+ // lt51-error at +2 {{use of undeclared identifier 'id'}}
+ // lt51-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(present, mapper(id): s)
+ // lt51-warning at +1 {{missing ':' after ) - ignoring}}
+ #pragma omp target update to(mapper(id) present: s)
+ // lt51-error at +2 {{use of undeclared identifier 'present'}}
+ // lt51-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(present mapper(id): s)
+
+ // Check parsing with unnecessary commas.
+ // lt51-warning at +1 {{missing ':' after ) - ignoring}}
+ #pragma omp target update to(mapper(id),: s)
+ // lt51-error at +3 {{use of undeclared identifier 'present'}}
+ // lt51-error at +2 {{expected expression}}
+ // lt51-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(present , : s)
+ // ge51-warning at +2 {{missing ':' after motion modifier - ignoring}}
+ // lt51-warning at +1 {{missing ':' after ) - ignoring}}
+ #pragma omp target update to(mapper(id),,: s)
+ // ge51-warning at +5 {{missing ':' after motion modifier - ignoring}}
+ // lt51-error at +4 {{use of undeclared identifier 'present'}}
+ // lt51-error at +3 {{expected expression}}
+ // lt51-error at +2 {{expected expression}}
+ // lt51-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(present,,: s)
+ // lt51-warning at +1 {{missing ':' after ) - ignoring}}
+ #pragma omp target update to(mapper(id), present,: s)
+ // lt51-error at +4 {{use of undeclared identifier 'present'}}
+ // lt51-error at +3 {{use of undeclared identifier 'id'}}
+ // lt51-error at +2 {{expected expression}}
+ // lt51-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+ #pragma omp target update to(present, mapper(id),: s)
+
#pragma omp target update from(m) allocate(m) // expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp target update'}}
{
foo();
}
int iarr[5][5];
-#pragma omp target update to(iarr[0:][1:2:-1]) // omp50-error {{section stride is evaluated to a non-positive value -1}} omp45-error {{expected ']'}} omp45-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+// ge50-error at +4 {{section stride is evaluated to a non-positive value -1}}
+// lt50-error at +3 {{expected ']'}}
+// lt50-note at +2 {{to match this '['}}
+// expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(iarr[0:][1:2:-1])
{}
-#pragma omp target update from(iarr[0:][1:2:-1]) // omp50-error {{section stride is evaluated to a non-positive value -1}} omp45-error {{expected ']'}} omp45-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+// ge50-error at +4 {{section stride is evaluated to a non-positive value -1}}
+// lt50-error at +3 {{expected ']'}}
+// lt50-note at +2 {{to match this '['}}
+// expected-error at +1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(iarr[0:][1:2:-1])
return tmain(argc, argv);
}
More information about the cfe-commits
mailing list