Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2024-11-24 06:10:12 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

[OpenMP] Code generation for the is_device_ptr clause

Browse Source 
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

llvm-svn: 276978
This commit is contained in:
Samuel Antao 2016-07-28 14:25:09 +00:00
parent cc10b85789
commit 6890b09634
9 changed files with 634 additions and 70 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

86
clang/include/clang/AST/OpenMPClause.h
View File

@ -4396,50 +4396,94 @@ public:
/// 'is_device_ptr' with the variables 'a' and 'b'. /// 'is_device_ptr' with the variables 'a' and 'b'.
/// ///
class OMPIsDevicePtrClause final class OMPIsDevicePtrClause final
: public OMPVarListClause<OMPIsDevicePtrClause>, : public OMPMappableExprListClause<OMPIsDevicePtrClause>,
private llvm::TrailingObjects<OMPIsDevicePtrClause, Expr *> { private llvm::TrailingObjects<
OMPIsDevicePtrClause, Expr *, ValueDecl *, unsigned,
OMPClauseMappableExprCommon::MappableComponent> {
friend TrailingObjects; friend TrailingObjects;
friend OMPVarListClause; friend OMPVarListClause;
friend OMPMappableExprListClause;
friend class OMPClauseReader; 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 StartLoc Starting location of the clause.
/// \param LParenLoc Location of '('.
/// \param EndLoc Ending location of the clause. /// \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, explicit OMPIsDevicePtrClause(SourceLocation StartLoc,
SourceLocation EndLoc, unsigned N) SourceLocation LParenLoc, SourceLocation EndLoc,
: OMPVarListClause<OMPIsDevicePtrClause>(OMPC_is_device_ptr, StartLoc, unsigned NumVars,
LParenLoc, EndLoc, N) {} unsigned NumUniqueDeclarations,
unsigned NumComponentLists,
unsigned NumComponents)
: OMPMappableExprListClause(OMPC_is_device_ptr, StartLoc, LParenLoc,
EndLoc, NumVars, NumUniqueDeclarations,
NumComponentLists, NumComponents) {}
/// Build an empty clause. /// 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) explicit OMPIsDevicePtrClause(unsigned NumVars,
: OMPVarListClause<OMPIsDevicePtrClause>( unsigned NumUniqueDeclarations,
OMPC_is_device_ptr, SourceLocation(), SourceLocation(), unsigned NumComponentLists,
SourceLocation(), N) {} unsigned NumComponents)
: OMPMappableExprListClause(OMPC_is_device_ptr, SourceLocation(),
SourceLocation(), SourceLocation(), NumVars,
NumUniqueDeclarations, NumComponentLists,
NumComponents) {}
public: public:
/// Creates clause with a list of variables \a VL. /// Creates clause with a list of variables \a Vars.
/// ///
/// \param C AST context. /// \param C AST context.
/// \param StartLoc Starting location of the clause. /// \param StartLoc Starting location of the clause.
/// \param LParenLoc Location of '('.
/// \param EndLoc Ending location of the clause. /// \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 * static OMPIsDevicePtrClause *
Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc, Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc,
SourceLocation EndLoc, ArrayRef<Expr *> VL); SourceLocation EndLoc, ArrayRef<Expr *> Vars,
/// Creates an empty clause with the place for \a N variables. ArrayRef<ValueDecl *> Declarations,
MappableExprComponentListsRef ComponentLists);
/// Creates an empty clause with the place for \a NumVars variables.
/// ///
/// \param C AST context. /// \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() { child_range children() {
return child_range(reinterpret_cast<Stmt **>(varlist_begin()), return child_range(reinterpret_cast<Stmt **>(varlist_begin()),

4
clang/include/clang/Basic/DiagnosticSemaKinds.td
View File

@ -8372,8 +8372,8 @@ def err_omp_schedule_nonmonotonic_ordered : Error<
"'schedule' clause with 'nonmonotonic' modifier cannot be specified if an 'ordered' clause is specified">; "'schedule' clause with 'nonmonotonic' modifier cannot be specified if an 'ordered' clause is specified">;
def err_omp_ordered_simd : Error< def err_omp_ordered_simd : Error<
"'ordered' clause with a parameter can not be specified in '#pragma omp %0' directive">; "'ordered' clause with a parameter can not be specified in '#pragma omp %0' directive">;
def err_omp_variable_in_map_and_dsa : Error< def err_omp_variable_in_given_clause_and_dsa : Error<
"%0 variable cannot be in a map clause in '#pragma omp %1' directive">; "%0 variable cannot be in a %1 clause in '#pragma omp %2' directive">;
def err_omp_param_or_this_in_clause : Error< def err_omp_param_or_this_in_clause : Error<
"expected reference to one of the parameters of function %0%select{| or 'this'}1">; "expected reference to one of the parameters of function %0%select{| or 'this'}1">;
def err_omp_expected_uniform_param : Error< def err_omp_expected_uniform_param : Error<

57
clang/lib/AST/OpenMPClause.cpp
View File

@ -794,20 +794,51 @@ OMPUseDevicePtrClause *OMPUseDevicePtrClause::CreateEmpty(
NumComponentLists, NumComponents); NumComponentLists, NumComponents);
} }
OMPIsDevicePtrClause *OMPIsDevicePtrClause::Create(const ASTContext &C, OMPIsDevicePtrClause *
SourceLocation StartLoc, OMPIsDevicePtrClause::Create(const ASTContext &C, SourceLocation StartLoc,
SourceLocation LParenLoc, SourceLocation LParenLoc, SourceLocation EndLoc,
SourceLocation EndLoc, ArrayRef<Expr *> Vars,
ArrayRef<Expr *> VL) { ArrayRef<ValueDecl *> Declarations,
void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size())); MappableExprComponentListsRef ComponentLists) {
OMPIsDevicePtrClause *Clause = unsigned NumVars = Vars.size();
new (Mem) OMPIsDevicePtrClause(StartLoc, LParenLoc, EndLoc, VL.size()); unsigned NumUniqueDeclarations =
Clause->setVarRefs(VL); 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; return Clause;
} }
OMPIsDevicePtrClause *OMPIsDevicePtrClause::CreateEmpty(const ASTContext &C, OMPIsDevicePtrClause *OMPIsDevicePtrClause::CreateEmpty(
unsigned N) { const ASTContext &C, unsigned NumVars, unsigned NumUniqueDeclarations,
void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N)); unsigned NumComponentLists, unsigned NumComponents) {
return new (Mem) OMPIsDevicePtrClause(N); void *Mem = C.Allocate(
totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
OMPClauseMappableExprCommon::MappableComponent>(
NumVars, NumUniqueDeclarations,
NumUniqueDeclarations + NumComponentLists, NumComponents));
return new (Mem) OMPIsDevicePtrClause(NumVars, NumUniqueDeclarations,
NumComponentLists, NumComponents);
} }

44
clang/lib/CodeGen/CGOpenMPRuntime.cpp
View File

@ -5022,6 +5022,13 @@ private:
/// \brief Set of all first private variables in the current directive. /// \brief Set of all first private variables in the current directive.
llvm::SmallPtrSet<const VarDecl *, 8> FirstPrivateDecls; 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 { llvm::Value *getExprTypeSize(const Expr *E) const {
auto ExprTy = E->getType().getCanonicalType(); auto ExprTy = E->getType().getCanonicalType();
@ -5418,6 +5425,10 @@ public:
for (const auto *D : C->varlists()) for (const auto *D : C->varlists())
FirstPrivateDecls.insert( FirstPrivateDecls.insert(
cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl())->getCanonicalDecl()); 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 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 /// \brief Generate the base pointers, section pointers, sizes and map types
/// associated to a given capture. /// associated to a given capture.
void generateInfoForCapture(const CapturedStmt::Capture *Cap, void generateInfoForCapture(const CapturedStmt::Capture *Cap,
llvm::Value *Arg,
MapBaseValuesArrayTy &BasePointers, MapBaseValuesArrayTy &BasePointers,
MapValuesArrayTy &Pointers, MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes, MapValuesArrayTy &Sizes,
@ -5585,14 +5597,38 @@ public:
Sizes.clear(); Sizes.clear();
Types.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 = const ValueDecl *VD =
Cap->capturesThis() Cap->capturesThis()
? nullptr ? nullptr
: cast<ValueDecl>(Cap->getCapturedVar()->getCanonicalDecl()); : cast<ValueDecl>(Cap->getCapturedVar()->getCanonicalDecl());
// We need to know when we generating information for the first component // If this declaration appears in a is_device_ptr clause we just have to
// associated with a capture, because the mapping flags depend on it. // pass the pointer by value. If it is a reference to a declaration, we just
bool IsFirstComponentList = true; // 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 *C : Directive.getClausesOfKind<OMPMapClause>())
for (auto L : C->decl_component_lists(VD)) { for (auto L : C->decl_component_lists(VD)) {
assert(L.first == VD && assert(L.first == VD &&
@ -5883,7 +5919,7 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
} else { } else {
// If we have any information in the map clause, we use it, otherwise we // If we have any information in the map clause, we use it, otherwise we
// just do a default mapping. // just do a default mapping.
MEHandler.generateInfoForCapture(CI, CurBasePointers, CurPointers, MEHandler.generateInfoForCapture(CI, *CV, CurBasePointers, CurPointers,
CurSizes, CurMapTypes); CurSizes, CurMapTypes);
if (CurBasePointers.empty()) if (CurBasePointers.empty())
MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers, MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers,

124
clang/lib/Sema/SemaOpenMP.cpp
View File

@ -72,8 +72,13 @@ private:
typedef llvm::DenseMap<ValueDecl *, Expr *> AlignedMapTy; typedef llvm::DenseMap<ValueDecl *, Expr *> AlignedMapTy;
typedef std::pair<unsigned, VarDecl *> LCDeclInfo; typedef std::pair<unsigned, VarDecl *> LCDeclInfo;
typedef llvm::DenseMap<ValueDecl *, LCDeclInfo> LoopControlVariablesMapTy; typedef llvm::DenseMap<ValueDecl *, LCDeclInfo> LoopControlVariablesMapTy;
typedef llvm::DenseMap< /// Struct that associates a component with the clause kind where they are
ValueDecl *, OMPClauseMappableExprCommon::MappableExprComponentLists> /// found.
struct MappedExprComponentTy {
OMPClauseMappableExprCommon::MappableExprComponentLists Components;
OpenMPClauseKind Kind = OMPC_unknown;
};
typedef llvm::DenseMap<ValueDecl *, MappedExprComponentTy>
MappedExprComponentsTy; MappedExprComponentsTy;
typedef llvm::StringMap<std::pair<OMPCriticalDirective *, llvm::APSInt>> typedef llvm::StringMap<std::pair<OMPCriticalDirective *, llvm::APSInt>>
CriticalsWithHintsTy; CriticalsWithHintsTy;
@ -327,8 +332,9 @@ public:
// if any issue is found. // if any issue is found.
bool checkMappableExprComponentListsForDecl( bool checkMappableExprComponentListsForDecl(
ValueDecl *VD, bool CurrentRegionOnly, ValueDecl *VD, bool CurrentRegionOnly,
const llvm::function_ref<bool( const llvm::function_ref<
OMPClauseMappableExprCommon::MappableExprComponentListRef)> &Check) { bool(OMPClauseMappableExprCommon::MappableExprComponentListRef,
OpenMPClauseKind)> &Check) {
auto SI = Stack.rbegin(); auto SI = Stack.rbegin();
auto SE = Stack.rend(); auto SE = Stack.rend();
@ -344,8 +350,8 @@ public:
for (; SI != SE; ++SI) { for (; SI != SE; ++SI) {
auto MI = SI->MappedExprComponents.find(VD); auto MI = SI->MappedExprComponents.find(VD);
if (MI != SI->MappedExprComponents.end()) if (MI != SI->MappedExprComponents.end())
for (auto &L : MI->second) for (auto &L : MI->second.Components)
if (Check(L)) if (Check(L, MI->second.Kind))
return true; return true;
} }
return false; return false;
@ -355,13 +361,15 @@ public:
// declaration and initialize it with the provided list of components. // declaration and initialize it with the provided list of components.
void addMappableExpressionComponents( void addMappableExpressionComponents(
ValueDecl *VD, ValueDecl *VD,
OMPClauseMappableExprCommon::MappableExprComponentListRef Components) { OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
OpenMPClauseKind WhereFoundClauseKind) {
assert(Stack.size() > 1 && assert(Stack.size() > 1 &&
"Not expecting to retrieve components from a empty stack!"); "Not expecting to retrieve components from a empty stack!");
auto &MEC = Stack.back().MappedExprComponents[VD]; auto &MEC = Stack.back().MappedExprComponents[VD];
// Create new entry and append the new components there. // Create new entry and append the new components there.
MEC.resize(MEC.size() + 1); MEC.Components.resize(MEC.Components.size() + 1);
MEC.back().append(Components.begin(), Components.end()); MEC.Components.back().append(Components.begin(), Components.end());
MEC.Kind = WhereFoundClauseKind;
} }
unsigned getNestingLevel() const { unsigned getNestingLevel() const {
@ -910,7 +918,13 @@ bool Sema::IsOpenMPCapturedByRef(ValueDecl *D, unsigned Level) {
DSAStack->checkMappableExprComponentListsForDecl( DSAStack->checkMappableExprComponentListsForDecl(
D, /*CurrentRegionOnly=*/true, D, /*CurrentRegionOnly=*/true,
[&](OMPClauseMappableExprCommon::MappableExprComponentListRef [&](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 EI = MapExprComponents.rbegin();
auto EE = MapExprComponents.rend(); auto EE = MapExprComponents.rend();
@ -8355,12 +8369,17 @@ OMPClause *Sema::ActOnOpenMPPrivateClause(ArrayRef<Expr *> VarList,
// A list item cannot appear in both a map clause and a data-sharing // A list item cannot appear in both a map clause and a data-sharing
// attribute clause on the same construct // attribute clause on the same construct
if (DSAStack->getCurrentDirective() == OMPD_target) { if (DSAStack->getCurrentDirective() == OMPD_target) {
OpenMPClauseKind ConflictKind;
if (DSAStack->checkMappableExprComponentListsForDecl( if (DSAStack->checkMappableExprComponentListsForDecl(
VD, /* CurrentRegionOnly = */ true, VD, /* CurrentRegionOnly = */ true,
[&](OMPClauseMappableExprCommon::MappableExprComponentListRef) [&](OMPClauseMappableExprCommon::MappableExprComponentListRef,
-> bool { return true; })) { OpenMPClauseKind WhereFoundClauseKind) -> bool {
Diag(ELoc, diag::err_omp_variable_in_map_and_dsa) ConflictKind = WhereFoundClauseKind;
return true;
})) {
Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
<< getOpenMPClauseName(OMPC_private) << getOpenMPClauseName(OMPC_private)
<< getOpenMPClauseName(ConflictKind)
<< getOpenMPDirectiveName(DSAStack->getCurrentDirective()); << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
ReportOriginalDSA(*this, DSAStack, D, DVar); ReportOriginalDSA(*this, DSAStack, D, DVar);
continue; continue;
@ -8606,12 +8625,17 @@ OMPClause *Sema::ActOnOpenMPFirstprivateClause(ArrayRef<Expr *> VarList,
// A list item cannot appear in both a map clause and a data-sharing // A list item cannot appear in both a map clause and a data-sharing
// attribute clause on the same construct // attribute clause on the same construct
if (CurrDir == OMPD_target) { if (CurrDir == OMPD_target) {
OpenMPClauseKind ConflictKind;
if (DSAStack->checkMappableExprComponentListsForDecl( if (DSAStack->checkMappableExprComponentListsForDecl(
VD, /* CurrentRegionOnly = */ true, VD, /* CurrentRegionOnly = */ true,
[&](OMPClauseMappableExprCommon::MappableExprComponentListRef) [&](OMPClauseMappableExprCommon::MappableExprComponentListRef,
-> bool { return true; })) { OpenMPClauseKind WhereFoundClauseKind) -> bool {
Diag(ELoc, diag::err_omp_variable_in_map_and_dsa) ConflictKind = WhereFoundClauseKind;
return true;
})) {
Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
<< getOpenMPClauseName(OMPC_firstprivate) << getOpenMPClauseName(OMPC_firstprivate)
<< getOpenMPClauseName(ConflictKind)
<< getOpenMPDirectiveName(DSAStack->getCurrentDirective()); << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
ReportOriginalDSA(*this, DSAStack, D, DVar); ReportOriginalDSA(*this, DSAStack, D, DVar);
continue; continue;
@ -10763,7 +10787,8 @@ static bool CheckMapConflicts(
bool FoundError = DSAS->checkMappableExprComponentListsForDecl( bool FoundError = DSAS->checkMappableExprComponentListsForDecl(
VD, CurrentRegionOnly, VD, CurrentRegionOnly,
[&](OMPClauseMappableExprCommon::MappableExprComponentListRef [&](OMPClauseMappableExprCommon::MappableExprComponentListRef
StackComponents) -> bool { StackComponents,
OpenMPClauseKind) -> bool {
assert(!StackComponents.empty() && assert(!StackComponents.empty() &&
"Map clause expression with no components!"); "Map clause expression with no components!");
@ -11121,8 +11146,9 @@ checkMappableExpressionList(Sema &SemaRef, DSAStackTy *DSAS,
if (DKind == OMPD_target && VD) { if (DKind == OMPD_target && VD) {
auto DVar = DSAS->getTopDSA(VD, false); auto DVar = DSAS->getTopDSA(VD, false);
if (isOpenMPPrivate(DVar.CKind)) { 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(DVar.CKind)
<< getOpenMPClauseName(OMPC_map)
<< getOpenMPDirectiveName(DSAS->getCurrentDirective()); << getOpenMPDirectiveName(DSAS->getCurrentDirective());
ReportOriginalDSA(SemaRef, DSAS, CurDeclaration, DVar); ReportOriginalDSA(SemaRef, DSAS, CurDeclaration, DVar);
continue; continue;
@ -11135,7 +11161,8 @@ checkMappableExpressionList(Sema &SemaRef, DSAStackTy *DSAS,
// Store the components in the stack so that they can be used to check // Store the components in the stack so that they can be used to check
// against other clauses later on. // 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 // Save the components and declaration to create the clause. For purposes of
// the clause creation, any component list that has has base 'this' uses // the clause creation, any component list that has has base 'this' uses
@ -11885,7 +11912,7 @@ OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
SourceLocation StartLoc, SourceLocation StartLoc,
SourceLocation LParenLoc, SourceLocation LParenLoc,
SourceLocation EndLoc) { SourceLocation EndLoc) {
SmallVector<Expr *, 8> Vars; MappableVarListInfo MVLI(VarList);
for (auto &RefExpr : VarList) { for (auto &RefExpr : VarList) {
assert(RefExpr && "NULL expr in OpenMP use_device_ptr clause."); assert(RefExpr && "NULL expr in OpenMP use_device_ptr clause.");
SourceLocation ELoc; SourceLocation ELoc;
@ -11894,7 +11921,7 @@ OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange); auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange);
if (Res.second) { if (Res.second) {
// It will be analyzed later. // It will be analyzed later.
Vars.push_back(RefExpr); MVLI.ProcessedVarList.push_back(RefExpr);
} }
ValueDecl *D = Res.first; ValueDecl *D = Res.first;
if (!D) if (!D)
@ -11908,12 +11935,59 @@ OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
<< 0 << RefExpr->getSourceRange(); << 0 << RefExpr->getSourceRange();
continue; 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 nullptr;
return OMPIsDevicePtrClause::Create(Context, StartLoc, LParenLoc, EndLoc, return OMPIsDevicePtrClause::Create(
Vars); Context, StartLoc, LParenLoc, EndLoc, MVLI.ProcessedVarList,
MVLI.VarBaseDeclarations, MVLI.VarComponents);
} }

46
clang/lib/Serialization/ASTReaderStmt.cpp
View File

@ -1941,10 +1941,16 @@ OMPClause *OMPClauseReader::readClause() {
NumLists, NumComponents); NumLists, NumComponents);
break; break;
} }
case OMPC_is_device_ptr: case OMPC_is_device_ptr: {
C = OMPIsDevicePtrClause::CreateEmpty(Context, Record[Idx++]); 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; break;
} }
}
Visit(C); Visit(C);
C->setLocStart(Reader->ReadSourceLocation(Record, Idx)); C->setLocStart(Reader->ReadSourceLocation(Record, Idx));
C->setLocEnd(Reader->ReadSourceLocation(Record, Idx)); C->setLocEnd(Reader->ReadSourceLocation(Record, Idx));
@ -2515,13 +2521,47 @@ void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
void OMPClauseReader::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) { void OMPClauseReader::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx)); 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; SmallVector<Expr *, 16> Vars;
Vars.reserve(NumVars); Vars.reserve(NumVars);
for (unsigned i = 0; i != NumVars; ++i) for (unsigned i = 0; i != NumVars; ++i)
Vars.push_back(Reader->Reader.ReadSubExpr()); Vars.push_back(Reader->Reader.ReadSubExpr());
C->setVarRefs(Vars); C->setVarRefs(Vars);
Vars.clear(); 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);
} }
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//

16
clang/lib/Serialization/ASTWriterStmt.cpp
View File

@ -2175,9 +2175,21 @@ void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
void OMPClauseWriter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) { void OMPClauseWriter::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
Record.push_back(C->varlist_size()); 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()); Record.AddSourceLocation(C->getLParenLoc());
for (auto *VE : C->varlists()) { for (auto *E : C->varlists())
Record.AddStmt(VE); 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());
} }
} }

293
clang/test/OpenMP/target_is_device_ptr_codegen.cpp Normal file
View File

@ -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

34
clang/test/OpenMP/target_is_device_ptr_messages.cpp
View File

@ -142,6 +142,7 @@ T tmain(T argc) {
T *&z = k; T *&z = k;
T aa[10]; T aa[10];
auto &raa = aa; 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 '(' 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( // 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 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; return 0;
} }
@ -194,6 +211,7 @@ int main(int argc, char **argv) {
int *&z = k; int *&z = k;
int aa[10]; int aa[10];
auto &raa = aa; 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 '(' 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( // 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 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}} return tmain<int, 3>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}}
} }