r276978 - [OpenMP] Code generation for the is_device_ptr clause
Samuel Antao via cfe-commits
cfe-commits at lists.llvm.org
Thu Jul 28 07:25:09 PDT 2016
Author: sfantao
Date: Thu Jul 28 09:25:09 2016
New Revision: 276978
URL: http://llvm.org/viewvc/llvm-project?rev=276978&view=rev
Log:
[OpenMP] Code generation for the is_device_ptr clause
Summary: This patch adds support for the is_device_ptr clause. It expands SEMA to use the mappable expression logic that can only be tested with code generation in place and check conflicts with other data sharing related clauses using the mappable expressions infrastructure.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: https://reviews.llvm.org/D22788
Added:
cfe/trunk/test/OpenMP/target_is_device_ptr_codegen.cpp
Modified:
cfe/trunk/include/clang/AST/OpenMPClause.h
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/AST/OpenMPClause.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/lib/Serialization/ASTReaderStmt.cpp
cfe/trunk/lib/Serialization/ASTWriterStmt.cpp
cfe/trunk/test/OpenMP/target_is_device_ptr_messages.cpp
Modified: cfe/trunk/include/clang/AST/OpenMPClause.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/OpenMPClause.h?rev=276978&r1=276977&r2=276978&view=diff
==============================================================================
--- cfe/trunk/include/clang/AST/OpenMPClause.h (original)
+++ cfe/trunk/include/clang/AST/OpenMPClause.h Thu Jul 28 09:25:09 2016
@@ -4396,50 +4396,94 @@ public:
/// 'is_device_ptr' with the variables 'a' and 'b'.
///
class OMPIsDevicePtrClause final
- : public OMPVarListClause<OMPIsDevicePtrClause>,
- private llvm::TrailingObjects<OMPIsDevicePtrClause, Expr *> {
+ : public OMPMappableExprListClause<OMPIsDevicePtrClause>,
+ private llvm::TrailingObjects<
+ OMPIsDevicePtrClause, Expr *, ValueDecl *, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent> {
friend TrailingObjects;
friend OMPVarListClause;
+ friend OMPMappableExprListClause;
friend class OMPClauseReader;
- /// Build clause with number of variables \a N.
+
+ /// Define the sizes of each trailing object array except the last one. This
+ /// is required for TrailingObjects to work properly.
+ size_t numTrailingObjects(OverloadToken<Expr *>) const {
+ return varlist_size();
+ }
+ size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
+ return getUniqueDeclarationsNum();
+ }
+ size_t numTrailingObjects(OverloadToken<unsigned>) const {
+ return getUniqueDeclarationsNum() + getTotalComponentListNum();
+ }
+ /// Build clause with number of variables \a NumVars.
///
/// \param StartLoc Starting location of the clause.
- /// \param LParenLoc Location of '('.
/// \param EndLoc Ending location of the clause.
- /// \param N Number of the variables in the clause.
- ///
- OMPIsDevicePtrClause(SourceLocation StartLoc, SourceLocation LParenLoc,
- SourceLocation EndLoc, unsigned N)
- : OMPVarListClause<OMPIsDevicePtrClause>(OMPC_is_device_ptr, StartLoc,
- LParenLoc, EndLoc, N) {}
+ /// \param NumVars Number of expressions listed in this clause.
+ /// \param NumUniqueDeclarations Number of unique base declarations in this
+ /// clause.
+ /// \param NumComponentLists Number of component lists in this clause.
+ /// \param NumComponents Total number of expression components in the clause.
+ ///
+ explicit OMPIsDevicePtrClause(SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc,
+ unsigned NumVars,
+ unsigned NumUniqueDeclarations,
+ unsigned NumComponentLists,
+ unsigned NumComponents)
+ : OMPMappableExprListClause(OMPC_is_device_ptr, StartLoc, LParenLoc,
+ EndLoc, NumVars, NumUniqueDeclarations,
+ NumComponentLists, NumComponents) {}
/// Build an empty clause.
///
- /// \param N Number of variables.
- ///
- explicit OMPIsDevicePtrClause(unsigned N)
- : OMPVarListClause<OMPIsDevicePtrClause>(
- OMPC_is_device_ptr, SourceLocation(), SourceLocation(),
- SourceLocation(), N) {}
+ /// \param NumVars Number of expressions listed in this clause.
+ /// \param NumUniqueDeclarations Number of unique base declarations in this
+ /// clause.
+ /// \param NumComponentLists Number of component lists in this clause.
+ /// \param NumComponents Total number of expression components in the clause.
+ ///
+ explicit OMPIsDevicePtrClause(unsigned NumVars,
+ unsigned NumUniqueDeclarations,
+ unsigned NumComponentLists,
+ unsigned NumComponents)
+ : OMPMappableExprListClause(OMPC_is_device_ptr, SourceLocation(),
+ SourceLocation(), SourceLocation(), NumVars,
+ NumUniqueDeclarations, NumComponentLists,
+ NumComponents) {}
public:
- /// Creates clause with a list of variables \a VL.
+ /// Creates clause with a list of variables \a Vars.
///
/// \param C AST context.
/// \param StartLoc Starting location of the clause.
- /// \param LParenLoc Location of '('.
/// \param EndLoc Ending location of the clause.
- /// \param VL List of references to the variables.
+ /// \param Vars The original expression used in the clause.
+ /// \param Declarations Declarations used in the clause.
+ /// \param ComponentLists Component lists used in the clause.
///
static OMPIsDevicePtrClause *
Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc,
- SourceLocation EndLoc, ArrayRef<Expr *> VL);
- /// Creates an empty clause with the place for \a N variables.
+ SourceLocation EndLoc, ArrayRef<Expr *> Vars,
+ ArrayRef<ValueDecl *> Declarations,
+ MappableExprComponentListsRef ComponentLists);
+
+ /// Creates an empty clause with the place for \a NumVars variables.
///
/// \param C AST context.
- /// \param N The number of variables.
- ///
- static OMPIsDevicePtrClause *CreateEmpty(const ASTContext &C, unsigned N);
+ /// \param NumVars Number of expressions listed in the clause.
+ /// \param NumUniqueDeclarations Number of unique base declarations in this
+ /// clause.
+ /// \param NumComponentLists Number of unique base declarations in this
+ /// clause.
+ /// \param NumComponents Total number of expression components in the clause.
+ ///
+ static OMPIsDevicePtrClause *CreateEmpty(const ASTContext &C,
+ unsigned NumVars,
+ unsigned NumUniqueDeclarations,
+ unsigned NumComponentLists,
+ unsigned NumComponents);
child_range children() {
return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=276978&r1=276977&r2=276978&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Thu Jul 28 09:25:09 2016
@@ -8372,8 +8372,8 @@ def err_omp_schedule_nonmonotonic_ordere
"'schedule' clause with 'nonmonotonic' modifier cannot be specified if an 'ordered' clause is specified">;
def err_omp_ordered_simd : Error<
"'ordered' clause with a parameter can not be specified in '#pragma omp %0' directive">;
-def err_omp_variable_in_map_and_dsa : Error<
- "%0 variable cannot be in a map clause in '#pragma omp %1' directive">;
+def err_omp_variable_in_given_clause_and_dsa : Error<
+ "%0 variable cannot be in a %1 clause in '#pragma omp %2' directive">;
def err_omp_param_or_this_in_clause : Error<
"expected reference to one of the parameters of function %0%select{| or 'this'}1">;
def err_omp_expected_uniform_param : Error<
Modified: cfe/trunk/lib/AST/OpenMPClause.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/OpenMPClause.cpp?rev=276978&r1=276977&r2=276978&view=diff
==============================================================================
--- cfe/trunk/lib/AST/OpenMPClause.cpp (original)
+++ cfe/trunk/lib/AST/OpenMPClause.cpp Thu Jul 28 09:25:09 2016
@@ -794,20 +794,51 @@ OMPUseDevicePtrClause *OMPUseDevicePtrCl
NumComponentLists, NumComponents);
}
-OMPIsDevicePtrClause *OMPIsDevicePtrClause::Create(const ASTContext &C,
- SourceLocation StartLoc,
- SourceLocation LParenLoc,
- SourceLocation EndLoc,
- ArrayRef<Expr *> VL) {
- void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
- OMPIsDevicePtrClause *Clause =
- new (Mem) OMPIsDevicePtrClause(StartLoc, LParenLoc, EndLoc, VL.size());
- Clause->setVarRefs(VL);
+OMPIsDevicePtrClause *
+OMPIsDevicePtrClause::Create(const ASTContext &C, SourceLocation StartLoc,
+ SourceLocation LParenLoc, SourceLocation EndLoc,
+ ArrayRef<Expr *> Vars,
+ ArrayRef<ValueDecl *> Declarations,
+ MappableExprComponentListsRef ComponentLists) {
+ unsigned NumVars = Vars.size();
+ unsigned NumUniqueDeclarations =
+ getUniqueDeclarationsTotalNumber(Declarations);
+ unsigned NumComponentLists = ComponentLists.size();
+ unsigned NumComponents = getComponentsTotalNumber(ComponentLists);
+
+ // We need to allocate:
+ // NumVars x Expr* - we have an original list expression for each clause list
+ // entry.
+ // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
+ // with each component list.
+ // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
+ // number of lists for each unique declaration and the size of each component
+ // list.
+ // NumComponents x MappableComponent - the total of all the components in all
+ // the lists.
+ void *Mem = C.Allocate(
+ totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent>(
+ NumVars, NumUniqueDeclarations,
+ NumUniqueDeclarations + NumComponentLists, NumComponents));
+
+ OMPIsDevicePtrClause *Clause = new (Mem) OMPIsDevicePtrClause(
+ StartLoc, LParenLoc, EndLoc, NumVars, NumUniqueDeclarations,
+ NumComponentLists, NumComponents);
+
+ Clause->setVarRefs(Vars);
+ Clause->setClauseInfo(Declarations, ComponentLists);
return Clause;
}
-OMPIsDevicePtrClause *OMPIsDevicePtrClause::CreateEmpty(const ASTContext &C,
- unsigned N) {
- void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));
- return new (Mem) OMPIsDevicePtrClause(N);
+OMPIsDevicePtrClause *OMPIsDevicePtrClause::CreateEmpty(
+ const ASTContext &C, unsigned NumVars, unsigned NumUniqueDeclarations,
+ unsigned NumComponentLists, unsigned NumComponents) {
+ void *Mem = C.Allocate(
+ totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
+ OMPClauseMappableExprCommon::MappableComponent>(
+ NumVars, NumUniqueDeclarations,
+ NumUniqueDeclarations + NumComponentLists, NumComponents));
+ return new (Mem) OMPIsDevicePtrClause(NumVars, NumUniqueDeclarations,
+ NumComponentLists, NumComponents);
}
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=276978&r1=276977&r2=276978&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Thu Jul 28 09:25:09 2016
@@ -5022,6 +5022,13 @@ private:
/// \brief Set of all first private variables in the current directive.
llvm::SmallPtrSet<const VarDecl *, 8> FirstPrivateDecls;
+ /// Map between device pointer declarations and their expression components.
+ /// The key value for declarations in 'this' is null.
+ llvm::DenseMap<
+ const ValueDecl *,
+ SmallVector<OMPClauseMappableExprCommon::MappableExprComponentListRef, 4>>
+ DevPointersMap;
+
llvm::Value *getExprTypeSize(const Expr *E) const {
auto ExprTy = E->getType().getCanonicalType();
@@ -5418,6 +5425,10 @@ public:
for (const auto *D : C->varlists())
FirstPrivateDecls.insert(
cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl())->getCanonicalDecl());
+ // Extract device pointer clause information.
+ for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
+ for (auto L : C->component_lists())
+ DevPointersMap[L.first].push_back(L.second);
}
/// \brief Generate all the base pointers, section pointers, sizes and map
@@ -5573,6 +5584,7 @@ public:
/// \brief Generate the base pointers, section pointers, sizes and map types
/// associated to a given capture.
void generateInfoForCapture(const CapturedStmt::Capture *Cap,
+ llvm::Value *Arg,
MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes,
@@ -5585,14 +5597,38 @@ public:
Sizes.clear();
Types.clear();
+ // We need to know when we generating information for the first component
+ // associated with a capture, because the mapping flags depend on it.
+ bool IsFirstComponentList = true;
+
const ValueDecl *VD =
Cap->capturesThis()
? nullptr
: cast<ValueDecl>(Cap->getCapturedVar()->getCanonicalDecl());
- // We need to know when we generating information for the first component
- // associated with a capture, because the mapping flags depend on it.
- bool IsFirstComponentList = true;
+ // If this declaration appears in a is_device_ptr clause we just have to
+ // pass the pointer by value. If it is a reference to a declaration, we just
+ // pass its value, otherwise, if it is a member expression, we need to map
+ // 'to' the field.
+ if (!VD) {
+ auto It = DevPointersMap.find(VD);
+ if (It != DevPointersMap.end()) {
+ for (auto L : It->second) {
+ generateInfoForComponentList(
+ /*MapType=*/OMPC_MAP_to, /*MapTypeModifier=*/OMPC_MAP_unknown, L,
+ BasePointers, Pointers, Sizes, Types, IsFirstComponentList);
+ IsFirstComponentList = false;
+ }
+ return;
+ }
+ } else if (DevPointersMap.count(VD)) {
+ BasePointers.push_back({Arg, VD});
+ Pointers.push_back(Arg);
+ Sizes.push_back(CGF.getTypeSize(CGF.getContext().VoidPtrTy));
+ Types.push_back(OMP_MAP_PRIVATE_VAL | OMP_MAP_FIRST_REF);
+ return;
+ }
+
for (auto *C : Directive.getClausesOfKind<OMPMapClause>())
for (auto L : C->decl_component_lists(VD)) {
assert(L.first == VD &&
@@ -5883,7 +5919,7 @@ void CGOpenMPRuntime::emitTargetCall(Cod
} else {
// If we have any information in the map clause, we use it, otherwise we
// just do a default mapping.
- MEHandler.generateInfoForCapture(CI, CurBasePointers, CurPointers,
+ MEHandler.generateInfoForCapture(CI, *CV, CurBasePointers, CurPointers,
CurSizes, CurMapTypes);
if (CurBasePointers.empty())
MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers,
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=276978&r1=276977&r2=276978&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Thu Jul 28 09:25:09 2016
@@ -72,8 +72,13 @@ private:
typedef llvm::DenseMap<ValueDecl *, Expr *> AlignedMapTy;
typedef std::pair<unsigned, VarDecl *> LCDeclInfo;
typedef llvm::DenseMap<ValueDecl *, LCDeclInfo> LoopControlVariablesMapTy;
- typedef llvm::DenseMap<
- ValueDecl *, OMPClauseMappableExprCommon::MappableExprComponentLists>
+ /// Struct that associates a component with the clause kind where they are
+ /// found.
+ struct MappedExprComponentTy {
+ OMPClauseMappableExprCommon::MappableExprComponentLists Components;
+ OpenMPClauseKind Kind = OMPC_unknown;
+ };
+ typedef llvm::DenseMap<ValueDecl *, MappedExprComponentTy>
MappedExprComponentsTy;
typedef llvm::StringMap<std::pair<OMPCriticalDirective *, llvm::APSInt>>
CriticalsWithHintsTy;
@@ -327,8 +332,9 @@ public:
// if any issue is found.
bool checkMappableExprComponentListsForDecl(
ValueDecl *VD, bool CurrentRegionOnly,
- const llvm::function_ref<bool(
- OMPClauseMappableExprCommon::MappableExprComponentListRef)> &Check) {
+ const llvm::function_ref<
+ bool(OMPClauseMappableExprCommon::MappableExprComponentListRef,
+ OpenMPClauseKind)> &Check) {
auto SI = Stack.rbegin();
auto SE = Stack.rend();
@@ -344,8 +350,8 @@ public:
for (; SI != SE; ++SI) {
auto MI = SI->MappedExprComponents.find(VD);
if (MI != SI->MappedExprComponents.end())
- for (auto &L : MI->second)
- if (Check(L))
+ for (auto &L : MI->second.Components)
+ if (Check(L, MI->second.Kind))
return true;
}
return false;
@@ -355,13 +361,15 @@ public:
// declaration and initialize it with the provided list of components.
void addMappableExpressionComponents(
ValueDecl *VD,
- OMPClauseMappableExprCommon::MappableExprComponentListRef Components) {
+ OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
+ OpenMPClauseKind WhereFoundClauseKind) {
assert(Stack.size() > 1 &&
"Not expecting to retrieve components from a empty stack!");
auto &MEC = Stack.back().MappedExprComponents[VD];
// Create new entry and append the new components there.
- MEC.resize(MEC.size() + 1);
- MEC.back().append(Components.begin(), Components.end());
+ MEC.Components.resize(MEC.Components.size() + 1);
+ MEC.Components.back().append(Components.begin(), Components.end());
+ MEC.Kind = WhereFoundClauseKind;
}
unsigned getNestingLevel() const {
@@ -910,7 +918,13 @@ bool Sema::IsOpenMPCapturedByRef(ValueDe
DSAStack->checkMappableExprComponentListsForDecl(
D, /*CurrentRegionOnly=*/true,
[&](OMPClauseMappableExprCommon::MappableExprComponentListRef
- MapExprComponents) {
+ MapExprComponents,
+ OpenMPClauseKind WhereFoundClauseKind) {
+ // Only the map clause information influences how a variable is
+ // captured. E.g. is_device_ptr does not require changing the default
+ // behaviour.
+ if (WhereFoundClauseKind != OMPC_map)
+ return false;
auto EI = MapExprComponents.rbegin();
auto EE = MapExprComponents.rend();
@@ -8355,12 +8369,17 @@ OMPClause *Sema::ActOnOpenMPPrivateClaus
// A list item cannot appear in both a map clause and a data-sharing
// attribute clause on the same construct
if (DSAStack->getCurrentDirective() == OMPD_target) {
+ OpenMPClauseKind ConflictKind;
if (DSAStack->checkMappableExprComponentListsForDecl(
VD, /* CurrentRegionOnly = */ true,
- [&](OMPClauseMappableExprCommon::MappableExprComponentListRef)
- -> bool { return true; })) {
- Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
+ [&](OMPClauseMappableExprCommon::MappableExprComponentListRef,
+ OpenMPClauseKind WhereFoundClauseKind) -> bool {
+ ConflictKind = WhereFoundClauseKind;
+ return true;
+ })) {
+ Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
<< getOpenMPClauseName(OMPC_private)
+ << getOpenMPClauseName(ConflictKind)
<< getOpenMPDirectiveName(DSAStack->getCurrentDirective());
ReportOriginalDSA(*this, DSAStack, D, DVar);
continue;
@@ -8606,12 +8625,17 @@ OMPClause *Sema::ActOnOpenMPFirstprivate
// A list item cannot appear in both a map clause and a data-sharing
// attribute clause on the same construct
if (CurrDir == OMPD_target) {
+ OpenMPClauseKind ConflictKind;
if (DSAStack->checkMappableExprComponentListsForDecl(
VD, /* CurrentRegionOnly = */ true,
- [&](OMPClauseMappableExprCommon::MappableExprComponentListRef)
- -> bool { return true; })) {
- Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
+ [&](OMPClauseMappableExprCommon::MappableExprComponentListRef,
+ OpenMPClauseKind WhereFoundClauseKind) -> bool {
+ ConflictKind = WhereFoundClauseKind;
+ return true;
+ })) {
+ Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
<< getOpenMPClauseName(OMPC_firstprivate)
+ << getOpenMPClauseName(ConflictKind)
<< getOpenMPDirectiveName(DSAStack->getCurrentDirective());
ReportOriginalDSA(*this, DSAStack, D, DVar);
continue;
@@ -10763,7 +10787,8 @@ static bool CheckMapConflicts(
bool FoundError = DSAS->checkMappableExprComponentListsForDecl(
VD, CurrentRegionOnly,
[&](OMPClauseMappableExprCommon::MappableExprComponentListRef
- StackComponents) -> bool {
+ StackComponents,
+ OpenMPClauseKind) -> bool {
assert(!StackComponents.empty() &&
"Map clause expression with no components!");
@@ -11121,8 +11146,9 @@ checkMappableExpressionList(Sema &SemaRe
if (DKind == OMPD_target && VD) {
auto DVar = DSAS->getTopDSA(VD, false);
if (isOpenMPPrivate(DVar.CKind)) {
- SemaRef.Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
+ SemaRef.Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
<< getOpenMPClauseName(DVar.CKind)
+ << getOpenMPClauseName(OMPC_map)
<< getOpenMPDirectiveName(DSAS->getCurrentDirective());
ReportOriginalDSA(SemaRef, DSAS, CurDeclaration, DVar);
continue;
@@ -11135,7 +11161,8 @@ checkMappableExpressionList(Sema &SemaRe
// Store the components in the stack so that they can be used to check
// against other clauses later on.
- DSAS->addMappableExpressionComponents(CurDeclaration, CurComponents);
+ DSAS->addMappableExpressionComponents(CurDeclaration, CurComponents,
+ /*WhereFoundClauseKind=*/OMPC_map);
// Save the components and declaration to create the clause. For purposes of
// the clause creation, any component list that has has base 'this' uses
@@ -11885,7 +11912,7 @@ OMPClause *Sema::ActOnOpenMPIsDevicePtrC
SourceLocation StartLoc,
SourceLocation LParenLoc,
SourceLocation EndLoc) {
- SmallVector<Expr *, 8> Vars;
+ MappableVarListInfo MVLI(VarList);
for (auto &RefExpr : VarList) {
assert(RefExpr && "NULL expr in OpenMP use_device_ptr clause.");
SourceLocation ELoc;
@@ -11894,7 +11921,7 @@ OMPClause *Sema::ActOnOpenMPIsDevicePtrC
auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange);
if (Res.second) {
// It will be analyzed later.
- Vars.push_back(RefExpr);
+ MVLI.ProcessedVarList.push_back(RefExpr);
}
ValueDecl *D = Res.first;
if (!D)
@@ -11908,12 +11935,59 @@ OMPClause *Sema::ActOnOpenMPIsDevicePtrC
<< 0 << RefExpr->getSourceRange();
continue;
}
- Vars.push_back(RefExpr->IgnoreParens());
+
+ // Check if the declaration in the clause does not show up in any data
+ // sharing attribute.
+ auto DVar = DSAStack->getTopDSA(D, false);
+ if (isOpenMPPrivate(DVar.CKind)) {
+ Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
+ << getOpenMPClauseName(DVar.CKind)
+ << getOpenMPClauseName(OMPC_is_device_ptr)
+ << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
+ ReportOriginalDSA(*this, DSAStack, D, DVar);
+ continue;
+ }
+
+ Expr *ConflictExpr;
+ if (DSAStack->checkMappableExprComponentListsForDecl(
+ D, /* CurrentRegionOnly = */ true,
+ [&ConflictExpr](
+ OMPClauseMappableExprCommon::MappableExprComponentListRef R,
+ OpenMPClauseKind) -> bool {
+ ConflictExpr = R.front().getAssociatedExpression();
+ return true;
+ })) {
+ Diag(ELoc, diag::err_omp_map_shared_storage) << RefExpr->getSourceRange();
+ Diag(ConflictExpr->getExprLoc(), diag::note_used_here)
+ << ConflictExpr->getSourceRange();
+ continue;
+ }
+
+ // Store the components in the stack so that they can be used to check
+ // against other clauses later on.
+ OMPClauseMappableExprCommon::MappableComponent MC(SimpleRefExpr, D);
+ DSAStack->addMappableExpressionComponents(
+ D, MC, /*WhereFoundClauseKind=*/OMPC_is_device_ptr);
+
+ // Record the expression we've just processed.
+ MVLI.ProcessedVarList.push_back(SimpleRefExpr);
+
+ // Create a mappable component for the list item. List items in this clause
+ // only need a component. We use a null declaration to signal fields in
+ // 'this'.
+ assert((isa<DeclRefExpr>(SimpleRefExpr) ||
+ isa<CXXThisExpr>(cast<MemberExpr>(SimpleRefExpr)->getBase())) &&
+ "Unexpected device pointer expression!");
+ MVLI.VarBaseDeclarations.push_back(
+ isa<DeclRefExpr>(SimpleRefExpr) ? D : nullptr);
+ MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
+ MVLI.VarComponents.back().push_back(MC);
}
- if (Vars.empty())
+ if (MVLI.ProcessedVarList.empty())
return nullptr;
- return OMPIsDevicePtrClause::Create(Context, StartLoc, LParenLoc, EndLoc,
- Vars);
+ return OMPIsDevicePtrClause::Create(
+ Context, StartLoc, LParenLoc, EndLoc, MVLI.ProcessedVarList,
+ MVLI.VarBaseDeclarations, MVLI.VarComponents);
}
Modified: cfe/trunk/lib/Serialization/ASTReaderStmt.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTReaderStmt.cpp?rev=276978&r1=276977&r2=276978&view=diff
==============================================================================
--- cfe/trunk/lib/Serialization/ASTReaderStmt.cpp (original)
+++ cfe/trunk/lib/Serialization/ASTReaderStmt.cpp Thu Jul 28 09:25:09 2016
@@ -1941,10 +1941,16 @@ OMPClause *OMPClauseReader::readClause()
NumLists, NumComponents);
break;
}
- case OMPC_is_device_ptr:
- C = OMPIsDevicePtrClause::CreateEmpty(Context, Record[Idx++]);
+ case OMPC_is_device_ptr: {
+ unsigned NumVars = Record[Idx++];
+ unsigned NumDeclarations = Record[Idx++];
+ unsigned NumLists = Record[Idx++];
+ unsigned NumComponents = Record[Idx++];
+ C = OMPIsDevicePtrClause::CreateEmpty(Context, NumVars, NumDeclarations,
+ NumLists, NumComponents);
break;
}
+ }
Visit(C);
C->setLocStart(Reader->ReadSourceLocation(Record, Idx));
C->setLocEnd(Reader->ReadSourceLocation(Record, Idx));
@@ -2515,13 +2521,47 @@ void OMPClauseReader::VisitOMPUseDeviceP
void OMPClauseReader::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
- unsigned NumVars = C->varlist_size();
+ auto NumVars = C->varlist_size();
+ auto UniqueDecls = C->getUniqueDeclarationsNum();
+ auto TotalLists = C->getTotalComponentListNum();
+ auto TotalComponents = C->getTotalComponentsNum();
+
SmallVector<Expr *, 16> Vars;
Vars.reserve(NumVars);
for (unsigned i = 0; i != NumVars; ++i)
Vars.push_back(Reader->Reader.ReadSubExpr());
C->setVarRefs(Vars);
Vars.clear();
+
+ SmallVector<ValueDecl *, 16> Decls;
+ Decls.reserve(UniqueDecls);
+ for (unsigned i = 0; i < UniqueDecls; ++i)
+ Decls.push_back(
+ Reader->Reader.ReadDeclAs<ValueDecl>(Reader->F, Record, Idx));
+ C->setUniqueDecls(Decls);
+
+ SmallVector<unsigned, 16> ListsPerDecl;
+ ListsPerDecl.reserve(UniqueDecls);
+ for (unsigned i = 0; i < UniqueDecls; ++i)
+ ListsPerDecl.push_back(Record[Idx++]);
+ C->setDeclNumLists(ListsPerDecl);
+
+ SmallVector<unsigned, 32> ListSizes;
+ ListSizes.reserve(TotalLists);
+ for (unsigned i = 0; i < TotalLists; ++i)
+ ListSizes.push_back(Record[Idx++]);
+ C->setComponentListSizes(ListSizes);
+
+ SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
+ Components.reserve(TotalComponents);
+ for (unsigned i = 0; i < TotalComponents; ++i) {
+ Expr *AssociatedExpr = Reader->Reader.ReadSubExpr();
+ ValueDecl *AssociatedDecl =
+ Reader->Reader.ReadDeclAs<ValueDecl>(Reader->F, Record, Idx);
+ Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
+ AssociatedExpr, AssociatedDecl));
+ }
+ C->setComponents(Components, ListSizes);
}
//===----------------------------------------------------------------------===//
Modified: cfe/trunk/lib/Serialization/ASTWriterStmt.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTWriterStmt.cpp?rev=276978&r1=276977&r2=276978&view=diff
==============================================================================
--- cfe/trunk/lib/Serialization/ASTWriterStmt.cpp (original)
+++ cfe/trunk/lib/Serialization/ASTWriterStmt.cpp Thu Jul 28 09:25:09 2016
@@ -2175,9 +2175,21 @@ void OMPClauseWriter::VisitOMPUseDeviceP
void OMPClauseWriter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
Record.push_back(C->varlist_size());
+ Record.push_back(C->getUniqueDeclarationsNum());
+ Record.push_back(C->getTotalComponentListNum());
+ Record.push_back(C->getTotalComponentsNum());
Record.AddSourceLocation(C->getLParenLoc());
- for (auto *VE : C->varlists()) {
- Record.AddStmt(VE);
+ for (auto *E : C->varlists())
+ Record.AddStmt(E);
+ for (auto *D : C->all_decls())
+ Record.AddDeclRef(D);
+ for (auto N : C->all_num_lists())
+ Record.push_back(N);
+ for (auto N : C->all_lists_sizes())
+ Record.push_back(N);
+ for (auto &M : C->all_components()) {
+ Record.AddStmt(M.getAssociatedExpression());
+ Record.AddDeclRef(M.getAssociatedDeclaration());
}
}
Added: cfe/trunk/test/OpenMP/target_is_device_ptr_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_is_device_ptr_codegen.cpp?rev=276978&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/target_is_device_ptr_codegen.cpp (added)
+++ cfe/trunk/test/OpenMP/target_is_device_ptr_codegen.cpp Thu Jul 28 09:25:09 2016
@@ -0,0 +1,293 @@
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
+// RUN: %clang_cc1 -DCK1 -fopenmp -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-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 CK1 --check-prefix CK1-64
+// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
+// RUN: %clang_cc1 -DCK1 -fopenmp -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-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 CK1 --check-prefix CK1-32
+#ifdef CK1
+
+double *g;
+
+// CK1: @g = global double*
+// CK1: [[SIZES00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]
+// CK1: [[TYPES00:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+
+// CK1: [[SIZES01:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
+// CK1: [[TYPES01:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+
+// CK1: [[SIZES02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
+// CK1: [[TYPES02:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+
+// CK1: [[SIZES03:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
+// CK1: [[TYPES03:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+
+// CK1: [[SIZES04:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
+// CK1: [[TYPES04:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+
+// CK1: [[SIZES05:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]
+// CK1: [[TYPES05:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+
+// CK1: [[SIZES06:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}]
+// CK1: [[TYPES06:@.+]] = {{.+}}constant [2 x i32] [i32 288, i32 288]
+
+// CK1-LABEL: @_Z3foo
+template<typename T>
+void foo(float *&lr, T *&tr) {
+ float *l;
+ T *t;
+
+ // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES00]]{{.+}}, {{.+}}[[TYPES00]]{{.+}})
+ // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+ // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
+ // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
+ // CK1-DAG: [[VALBP]] = bitcast double* [[VAL:%.+]] to i8*
+ // CK1-DAG: [[VALP]] = bitcast double* [[VAL]] to i8*
+ // CK1-DAG: [[VAL]] = load double*, double** [[ADDR:@g]],
+
+ // CK1: call void [[KERNEL:@.+]](double* [[VAL]])
+ #pragma omp target is_device_ptr(g)
+ {
+ ++g;
+ }
+
+ // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES01]]{{.+}}, {{.+}}[[TYPES01]]{{.+}})
+ // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+ // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
+ // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
+ // CK1-DAG: [[VALBP]] = bitcast float* [[VAL:%.+]] to i8*
+ // CK1-DAG: [[VALP]] = bitcast float* [[VAL]] to i8*
+ // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
+
+ // CK1: call void [[KERNEL:@.+]](float* [[VAL]])
+ #pragma omp target is_device_ptr(l)
+ {
+ ++l;
+ }
+
+ // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES02]]{{.+}}, {{.+}}[[TYPES02]]{{.+}})
+ // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+ // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
+ // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
+ // CK1-DAG: [[VALBP]] = bitcast i32* [[VAL:%.+]] to i8*
+ // CK1-DAG: [[VALP]] = bitcast i32* [[VAL]] to i8*
+ // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
+
+ // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
+ #pragma omp target is_device_ptr(t)
+ {
+ ++t;
+ }
+
+ // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES03]]{{.+}}, {{.+}}[[TYPES03]]{{.+}})
+ // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+ // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
+ // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
+ // CK1-DAG: [[VALBP]] = bitcast float* [[VAL:%.+]] to i8*
+ // CK1-DAG: [[VALP]] = bitcast float* [[VAL]] to i8*
+ // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]],
+ // CK1-DAG: [[ADDR]] = load float**, float*** [[ADDR2:%.+]],
+
+ // CK1: call void [[KERNEL:@.+]](float* [[VAL]])
+ #pragma omp target is_device_ptr(lr)
+ {
+ ++lr;
+ }
+
+ // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES04]]{{.+}}, {{.+}}[[TYPES04]]{{.+}})
+ // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+ // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
+ // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
+ // CK1-DAG: [[VALBP]] = bitcast i32* [[VAL:%.+]] to i8*
+ // CK1-DAG: [[VALP]] = bitcast i32* [[VAL]] to i8*
+ // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
+ // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
+
+ // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
+ #pragma omp target is_device_ptr(tr)
+ {
+ ++tr;
+ }
+
+ // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES05]]{{.+}}, {{.+}}[[TYPES05]]{{.+}})
+ // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+ // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
+ // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
+ // CK1-DAG: [[VALBP]] = bitcast i32* [[VAL:%.+]] to i8*
+ // CK1-DAG: [[VALP]] = bitcast i32* [[VAL]] to i8*
+ // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
+ // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
+
+ // CK1: call void [[KERNEL:@.+]](i32* [[VAL]])
+ #pragma omp target is_device_ptr(tr,lr)
+ {
+ ++tr;
+ }
+
+ // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES06]]{{.+}}, {{.+}}[[TYPES06]]{{.+}})
+ // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+ // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+ // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+ // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
+ // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
+ // CK1-DAG: [[VALBP]] = bitcast i32* [[VAL:%.+]] to i8*
+ // CK1-DAG: [[VALP]] = bitcast i32* [[VAL]] to i8*
+ // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
+ // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
+
+ // CK1-DAG: [[_BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
+ // CK1-DAG: [[_P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
+ // CK1-DAG: store i8* [[_VALBP:%.+]], i8** [[_BP1]],
+ // CK1-DAG: store i8* [[_VALP:%.+]], i8** [[_P1]],
+ // CK1-DAG: [[_VALBP]] = bitcast float* [[_VAL:%.+]] to i8*
+ // CK1-DAG: [[_VALP]] = bitcast float* [[_VAL]] to i8*
+ // CK1-DAG: [[_VAL]] = load float*, float** [[_ADDR:%.+]],
+ // CK1-DAG: [[_ADDR]] = load float**, float*** [[_ADDR2:%.+]],
+
+ // CK1: call void [[KERNEL:@.+]](i32* [[VAL]], float* [[_VAL]])
+ #pragma omp target is_device_ptr(tr,lr)
+ {
+ ++tr,++lr;
+ }
+}
+
+void bar(float *&a, int *&b) {
+ foo<int>(a,b);
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
+// RUN: %clang_cc1 -DCK2 -fopenmp -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-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 CK2 --check-prefix CK2-64
+// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
+// RUN: %clang_cc1 -DCK2 -fopenmp -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-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 CK2 --check-prefix CK2-32
+#ifdef CK2
+
+// CK2: [[ST:%.+]] = type { double*, double** }
+
+// CK2: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 33]
+
+// CK2: [[SIZE01:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}]
+// CK2: [[MTYPE01:@.+]] = {{.+}}constant [2 x i32] [i32 32, i32 17]
+
+// CK2: [[SIZE02:@.+]] = {{.+}}constant [3 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}, i[[sz]] {{8|4}}]
+// CK2: [[MTYPE02:@.+]] = {{.+}}constant [3 x i32] [i32 33, i32 0, i32 17]
+
+template <typename T>
+struct ST {
+ T *a;
+ double *&b;
+ ST(double *&b) : a(0), b(b) {}
+
+ // CK2-LABEL: @{{.*}}foo{{.*}}
+ void foo(double *&arg) {
+ int *la = 0;
+
+ // CK2-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK2-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+ // CK2-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+ // CK2-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
+ // CK2-DAG: [[CPVAL0]] = bitcast double** [[SEC0:%.+]] to i8*
+ // CK2-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0
+ #pragma omp target is_device_ptr(a)
+ {
+ a++;
+ }
+
+ // CK2-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK2-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+ // CK2-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+ // CK2-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
+ // CK2-DAG: [[CPVAL0]] = bitcast double*** [[SEC0:%.+]] to i8*
+ // CK2-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+ // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK2-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+ // CK2-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+ // CK2-DAG: [[CBPVAL1]] = bitcast double*** [[SEC0]] to i8*
+ // CK2-DAG: [[CPVAL1]] = bitcast double** [[SEC1:%.+]] to i8*
+ // CK2-DAG: [[SEC1]] = load double**, double*** [[SEC0]]
+ #pragma omp target is_device_ptr(b)
+ {
+ b++;
+ }
+
+ // CK2-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE02]]{{.+}})
+ // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK2-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+ // CK2-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+ // CK2-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8*
+ // CK2-DAG: [[CPVAL0]] = bitcast double*** [[SEC0:%.+]] to i8*
+ // CK2-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
+
+ // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+ // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+ // CK2-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]]
+ // CK2-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]]
+ // CK2-DAG: [[CBPVAL1]] = bitcast double*** [[SEC0]] to i8*
+ // CK2-DAG: [[CPVAL1]] = bitcast double** [[SEC1:%.+]] to i8*
+ // CK2-DAG: [[SEC1]] = load double**, double*** [[SEC0]]
+
+ // CK2-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK2-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK2-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]]
+ // CK2-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]]
+ // CK2-DAG: [[CBPVAL2]] = bitcast [[ST]]* [[VAR2:%.+]] to i8*
+ // CK2-DAG: [[CPVAL2]] = bitcast double** [[SEC2:%.+]] to i8*
+ // CK2-DAG: [[SEC2]] = getelementptr {{.*}}[[ST]]* [[VAR2]], i{{.+}} 0, i{{.+}} 0
+ #pragma omp target is_device_ptr(a, b)
+ {
+ a++;
+ b++;
+ }
+ }
+};
+
+void bar(double *arg){
+ ST<double> A(arg);
+ A.foo(arg);
+ ++arg;
+}
+#endif
+#endif
Modified: cfe/trunk/test/OpenMP/target_is_device_ptr_messages.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_is_device_ptr_messages.cpp?rev=276978&r1=276977&r2=276978&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_is_device_ptr_messages.cpp (original)
+++ cfe/trunk/test/OpenMP/target_is_device_ptr_messages.cpp Thu Jul 28 09:25:09 2016
@@ -142,6 +142,7 @@ T tmain(T argc) {
T *&z = k;
T aa[10];
auto &raa = aa;
+ S6 *ps;
#pragma omp target is_device_ptr // expected-error {{expected '(' after 'is_device_ptr'}}
{}
#pragma omp target is_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
@@ -178,6 +179,22 @@ T tmain(T argc) {
{}
#pragma omp target is_device_ptr(da) // OK
{}
+#pragma omp target map(ps) is_device_ptr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target is_device_ptr(ps) map(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target map(ps->a) is_device_ptr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}}
+ {}
+#pragma omp target is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target' directive}}
+ {}
+#pragma omp target firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target' directive}} expected-note{{defined as firstprivate}}
+ {}
+#pragma omp target is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target' directive}}
+ {}
+#pragma omp target private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target' directive}} expected-note{{defined as private}}
+ {}
return 0;
}
@@ -194,6 +211,7 @@ int main(int argc, char **argv) {
int *&z = k;
int aa[10];
auto &raa = aa;
+ S6 *ps;
#pragma omp target is_device_ptr // expected-error {{expected '(' after 'is_device_ptr'}}
{}
#pragma omp target is_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
@@ -230,5 +248,21 @@ int main(int argc, char **argv) {
{}
#pragma omp target is_device_ptr(da) // OK
{}
+#pragma omp target map(ps) is_device_ptr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target is_device_ptr(ps) map(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target map(ps->a) is_device_ptr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+ {}
+#pragma omp target is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}}
+ {}
+#pragma omp target is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target' directive}}
+ {}
+#pragma omp target firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target' directive}} expected-note{{defined as firstprivate}}
+ {}
+#pragma omp target is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target' directive}}
+ {}
+#pragma omp target private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target' directive}} expected-note{{defined as private}}
+ {}
return tmain<int, 3>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}}
}
More information about the cfe-commits
mailing list