Index: include/clang/AST/OpenMPClause.h =================================================================== --- include/clang/AST/OpenMPClause.h +++ include/clang/AST/OpenMPClause.h @@ -4396,50 +4396,94 @@ /// 'is_device_ptr' with the variables 'a' and 'b'. /// class OMPIsDevicePtrClause final - : public OMPVarListClause, - private llvm::TrailingObjects { + : public OMPMappableExprListClause, + 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) const { + return varlist_size(); + } + size_t numTrailingObjects(OverloadToken) const { + return getUniqueDeclarationsNum(); + } + size_t numTrailingObjects(OverloadToken) 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(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( - 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 VL); - /// Creates an empty clause with the place for \a N variables. + SourceLocation EndLoc, ArrayRef Vars, + ArrayRef 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(varlist_begin()), Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -8372,8 +8372,8 @@ "'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< Index: lib/AST/OpenMPClause.cpp =================================================================== --- lib/AST/OpenMPClause.cpp +++ lib/AST/OpenMPClause.cpp @@ -794,20 +794,51 @@ NumComponentLists, NumComponents); } -OMPIsDevicePtrClause *OMPIsDevicePtrClause::Create(const ASTContext &C, - SourceLocation StartLoc, - SourceLocation LParenLoc, - SourceLocation EndLoc, - ArrayRef VL) { - void *Mem = C.Allocate(totalSizeToAlloc(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 Vars, + ArrayRef 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( + 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(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( + NumVars, NumUniqueDeclarations, + NumUniqueDeclarations + NumComponentLists, NumComponents)); + return new (Mem) OMPIsDevicePtrClause(NumVars, NumUniqueDeclarations, + NumComponentLists, NumComponents); } Index: lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenMPRuntime.cpp +++ lib/CodeGen/CGOpenMPRuntime.cpp @@ -5022,6 +5022,13 @@ /// \brief Set of all first private variables in the current directive. llvm::SmallPtrSet FirstPrivateDecls; + /// Map between device pointer declarations and their expression components. + /// The key value for declarations in 'this' is null. + llvm::DenseMap< + const ValueDecl *, + SmallVector> + DevPointersMap; + llvm::Value *getExprTypeSize(const Expr *E) const { auto ExprTy = E->getType().getCanonicalType(); @@ -5418,6 +5425,10 @@ for (const auto *D : C->varlists()) FirstPrivateDecls.insert( cast(cast(D)->getDecl())->getCanonicalDecl()); + // Extract device pointer clause information. + for (const auto *C : Dir.getClausesOfKind()) + 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 @@ /// \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 @@ 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(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()) for (auto L : C->decl_component_lists(VD)) { assert(L.first == VD && @@ -5883,7 +5919,7 @@ } 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, Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -72,8 +72,13 @@ typedef llvm::DenseMap AlignedMapTy; typedef std::pair LCDeclInfo; typedef llvm::DenseMap 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 MappedExprComponentsTy; typedef llvm::StringMap> CriticalsWithHintsTy; @@ -327,8 +332,9 @@ // if any issue is found. bool checkMappableExprComponentListsForDecl( ValueDecl *VD, bool CurrentRegionOnly, - const llvm::function_ref &Check) { + const llvm::function_ref< + bool(OMPClauseMappableExprCommon::MappableExprComponentListRef, + OpenMPClauseKind)> &Check) { auto SI = Stack.rbegin(); auto SE = Stack.rend(); @@ -344,8 +350,8 @@ 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 @@ // 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 @@ 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 @@ // 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 @@ // 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 @@ 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 @@ 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 @@ // 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 @@ SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc) { - SmallVector Vars; + MappableVarListInfo MVLI(VarList); for (auto &RefExpr : VarList) { assert(RefExpr && "NULL expr in OpenMP use_device_ptr clause."); SourceLocation ELoc; @@ -11894,7 +11921,7 @@ 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 @@ << 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(SimpleRefExpr) || + isa(cast(SimpleRefExpr)->getBase())) && + "Unexpected device pointer expression!"); + MVLI.VarBaseDeclarations.push_back( + isa(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); } Index: lib/Serialization/ASTReaderStmt.cpp =================================================================== --- lib/Serialization/ASTReaderStmt.cpp +++ lib/Serialization/ASTReaderStmt.cpp @@ -1941,10 +1941,16 @@ 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::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 Vars; Vars.reserve(NumVars); for (unsigned i = 0; i != NumVars; ++i) Vars.push_back(Reader->Reader.ReadSubExpr()); C->setVarRefs(Vars); Vars.clear(); + + SmallVector Decls; + Decls.reserve(UniqueDecls); + for (unsigned i = 0; i < UniqueDecls; ++i) + Decls.push_back( + Reader->Reader.ReadDeclAs(Reader->F, Record, Idx)); + C->setUniqueDecls(Decls); + + SmallVector ListsPerDecl; + ListsPerDecl.reserve(UniqueDecls); + for (unsigned i = 0; i < UniqueDecls; ++i) + ListsPerDecl.push_back(Record[Idx++]); + C->setDeclNumLists(ListsPerDecl); + + SmallVector ListSizes; + ListSizes.reserve(TotalLists); + for (unsigned i = 0; i < TotalLists; ++i) + ListSizes.push_back(Record[Idx++]); + C->setComponentListSizes(ListSizes); + + SmallVector Components; + Components.reserve(TotalComponents); + for (unsigned i = 0; i < TotalComponents; ++i) { + Expr *AssociatedExpr = Reader->Reader.ReadSubExpr(); + ValueDecl *AssociatedDecl = + Reader->Reader.ReadDeclAs(Reader->F, Record, Idx); + Components.push_back(OMPClauseMappableExprCommon::MappableComponent( + AssociatedExpr, AssociatedDecl)); + } + C->setComponents(Components, ListSizes); } //===----------------------------------------------------------------------===// Index: lib/Serialization/ASTWriterStmt.cpp =================================================================== --- lib/Serialization/ASTWriterStmt.cpp +++ lib/Serialization/ASTWriterStmt.cpp @@ -2175,9 +2175,21 @@ 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()); } } Index: test/OpenMP/target_is_device_ptr_codegen.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_is_device_ptr_codegen.cpp @@ -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 +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(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 +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 A(arg); + A.foo(arg); + ++arg; +} +#endif +#endif Index: test/OpenMP/target_is_device_ptr_messages.cpp =================================================================== --- test/OpenMP/target_is_device_ptr_messages.cpp +++ test/OpenMP/target_is_device_ptr_messages.cpp @@ -142,6 +142,7 @@ 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 @@ {} #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 *&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 @@ {} #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(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} }