/// '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.
+ /// \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.
///
- OMPIsDevicePtrClause(SourceLocation StartLoc, SourceLocation LParenLoc,
- SourceLocation EndLoc, unsigned N)
- : OMPVarListClause<OMPIsDevicePtrClause>(OMPC_is_device_ptr, StartLoc,
- LParenLoc, EndLoc, N) {}
+ 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.
+ /// \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 N)
- : OMPVarListClause<OMPIsDevicePtrClause>(
- OMPC_is_device_ptr, SourceLocation(), SourceLocation(),
- SourceLocation(), N) {}
+ 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.
+ /// \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 N);
+ 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()),
"'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<
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);
}
/// \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();
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
/// \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,
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 &&
} 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,
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;
// 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();
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;
// 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 {
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();
// 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;
// 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;
bool FoundError = DSAS->checkMappableExprComponentListsForDecl(
VD, CurrentRegionOnly,
[&](OMPClauseMappableExprCommon::MappableExprComponentListRef
- StackComponents) -> bool {
+ StackComponents,
+ OpenMPClauseKind) -> bool {
assert(!StackComponents.empty() &&
"Map clause expression with no components!");
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;
// 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
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;
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)
<< 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);
}
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));
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);
}
//===----------------------------------------------------------------------===//
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());
}
}
--- /dev/null
+// 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
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}}
{}
#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;
}
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}}
{}
#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}}
}
projects / clang / commitdiff