mirror of
https://github.com/intel/llvm.git
synced 2026-01-13 11:02:04 +08:00
[OpenMP] Permit map with DSA on combined directive
For `map`, the following restriction changed in OpenMP 5.0: * OpenMP 4.5 [2.15.5.1, Restrictions]: "A list item cannot appear in both a map clause and a data-sharing attribute clause on the same construct. * OpenMP 5.0 [2.19.7.1, Restrictions]: "A list item cannot appear in both a map clause and a data-sharing attribute clause on the same construct unless the construct is a combined construct." This patch removes this restriction in the case of combined constructs and OpenMP 5.0, and it updates Sema not to capture a scalar by copy in the target region when `firstprivate` and `map` appear for that scalar on a combined target construct. This patch also adds a fixme to a test that now reveals that a diagnostic about loop iteration variables is dropped in the case of OpenMP 5.0. That bug exists regardless of this patch's changes. Reviewed By: ABataev, jdoerfert, hfinkel, kkwli0 Differential Revision: https://reviews.llvm.org/D65835 llvm-svn: 369619
This commit is contained in:
Joel E. Denny
@@ -756,13 +756,16 @@ public:
|
||||
unsigned short CapRegionKind;
|
||||
|
||||
unsigned short OpenMPLevel;
|
||||
unsigned short OpenMPCaptureLevel;
|
||||
|
||||
CapturedRegionScopeInfo(DiagnosticsEngine &Diag, Scope *S, CapturedDecl *CD,
|
||||
RecordDecl *RD, ImplicitParamDecl *Context,
|
||||
CapturedRegionKind K, unsigned OpenMPLevel)
|
||||
CapturedRegionKind K, unsigned OpenMPLevel,
|
||||
unsigned OpenMPCaptureLevel)
|
||||
: CapturingScopeInfo(Diag, ImpCap_CapturedRegion),
|
||||
TheCapturedDecl(CD), TheRecordDecl(RD), TheScope(S),
|
||||
ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel) {
|
||||
ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel),
|
||||
OpenMPCaptureLevel(OpenMPCaptureLevel) {
|
||||
Kind = SK_CapturedRegion;
|
||||
}
|
||||
|
||||
|
||||
@@ -1421,8 +1421,8 @@ public:
|
||||
void RecordParsingTemplateParameterDepth(unsigned Depth);
|
||||
|
||||
void PushCapturedRegionScope(Scope *RegionScope, CapturedDecl *CD,
|
||||
RecordDecl *RD,
|
||||
CapturedRegionKind K);
|
||||
RecordDecl *RD, CapturedRegionKind K,
|
||||
unsigned OpenMPCaptureLevel = 0);
|
||||
|
||||
/// Custom deleter to allow FunctionScopeInfos to be kept alive for a short
|
||||
/// time after they've been popped.
|
||||
@@ -3979,7 +3979,8 @@ public:
|
||||
typedef std::pair<StringRef, QualType> CapturedParamNameType;
|
||||
void ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
|
||||
CapturedRegionKind Kind,
|
||||
ArrayRef<CapturedParamNameType> Params);
|
||||
ArrayRef<CapturedParamNameType> Params,
|
||||
unsigned OpenMPCaptureLevel = 0);
|
||||
StmtResult ActOnCapturedRegionEnd(Stmt *S);
|
||||
void ActOnCapturedRegionError();
|
||||
RecordDecl *CreateCapturedStmtRecordDecl(CapturedDecl *&CD,
|
||||
@@ -9028,7 +9029,9 @@ public:
|
||||
/// reference.
|
||||
/// \param Level Relative level of nested OpenMP construct for that the check
|
||||
/// is performed.
|
||||
bool isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const;
|
||||
/// \param OpenMPCaptureLevel Capture level within an OpenMP construct.
|
||||
bool isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
|
||||
unsigned OpenMPCaptureLevel) const;
|
||||
|
||||
/// Check if the specified variable is used in one of the private
|
||||
/// clauses (private, firstprivate, lastprivate, reduction etc.) in OpenMP
|
||||
|
||||
@@ -2105,10 +2105,12 @@ IdentifierInfo *Sema::getFloat128Identifier() const {
|
||||
}
|
||||
|
||||
void Sema::PushCapturedRegionScope(Scope *S, CapturedDecl *CD, RecordDecl *RD,
|
||||
CapturedRegionKind K) {
|
||||
CapturingScopeInfo *CSI = new CapturedRegionScopeInfo(
|
||||
CapturedRegionKind K,
|
||||
unsigned OpenMPCaptureLevel) {
|
||||
auto *CSI = new CapturedRegionScopeInfo(
|
||||
getDiagnostics(), S, CD, RD, CD->getContextParam(), K,
|
||||
(getLangOpts().OpenMP && K == CR_OpenMP) ? getOpenMPNestingLevel() : 0);
|
||||
(getLangOpts().OpenMP && K == CR_OpenMP) ? getOpenMPNestingLevel() : 0,
|
||||
OpenMPCaptureLevel);
|
||||
CSI->ReturnType = Context.VoidTy;
|
||||
FunctionScopes.push_back(CSI);
|
||||
}
|
||||
|
||||
@@ -15660,7 +15660,8 @@ static bool captureInCapturedRegion(CapturedRegionScopeInfo *RSI,
|
||||
if (HasConst)
|
||||
DeclRefType.addConst();
|
||||
}
|
||||
ByRef = S.isOpenMPCapturedByRef(Var, RSI->OpenMPLevel);
|
||||
ByRef = S.isOpenMPCapturedByRef(Var, RSI->OpenMPLevel,
|
||||
RSI->OpenMPCaptureLevel);
|
||||
}
|
||||
|
||||
if (ByRef)
|
||||
|
||||
@@ -522,6 +522,13 @@ public:
|
||||
assert(!isStackEmpty() && "No directive at specified level.");
|
||||
return getStackElemAtLevel(Level).Directive;
|
||||
}
|
||||
/// Returns the capture region at the specified level.
|
||||
OpenMPDirectiveKind getCaptureRegion(unsigned Level,
|
||||
unsigned OpenMPCaptureLevel) const {
|
||||
SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
|
||||
getOpenMPCaptureRegions(CaptureRegions, getDirective(Level));
|
||||
return CaptureRegions[OpenMPCaptureLevel];
|
||||
}
|
||||
/// Returns parent directive.
|
||||
OpenMPDirectiveKind getParentDirective() const {
|
||||
const SharingMapTy *Parent = getSecondOnStackOrNull();
|
||||
@@ -1340,7 +1347,8 @@ const DSAStackTy::DSAVarData DSAStackTy::getTopDSA(ValueDecl *D,
|
||||
}
|
||||
const_iterator End = end();
|
||||
if (!SemaRef.isOpenMPCapturedByRef(
|
||||
D, std::distance(ParentIterTarget, End))) {
|
||||
D, std::distance(ParentIterTarget, End),
|
||||
/*OpenMPCaptureLevel=*/0)) {
|
||||
DVar.RefExpr =
|
||||
buildDeclRefExpr(SemaRef, VD, D->getType().getNonReferenceType(),
|
||||
IterTarget->ConstructLoc);
|
||||
@@ -1611,7 +1619,8 @@ void Sema::checkOpenMPDeviceExpr(const Expr *E) {
|
||||
<< Context.getTargetInfo().getTriple().str() << E->getSourceRange();
|
||||
}
|
||||
|
||||
bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const {
|
||||
bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
|
||||
unsigned OpenMPCaptureLevel) const {
|
||||
assert(LangOpts.OpenMP && "OpenMP is not allowed");
|
||||
|
||||
ASTContext &Ctx = getASTContext();
|
||||
@@ -1621,6 +1630,7 @@ bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const {
|
||||
D = cast<ValueDecl>(D->getCanonicalDecl());
|
||||
QualType Ty = D->getType();
|
||||
|
||||
bool IsVariableUsedInMapClause = false;
|
||||
if (DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective, Level)) {
|
||||
// This table summarizes how a given variable should be passed to the device
|
||||
// given its type and the clauses where it appears. This table is based on
|
||||
@@ -1682,7 +1692,6 @@ bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const {
|
||||
// Locate map clauses and see if the variable being captured is referred to
|
||||
// in any of those clauses. Here we only care about variables, not fields,
|
||||
// because fields are part of aggregates.
|
||||
bool IsVariableUsedInMapClause = false;
|
||||
bool IsVariableAssociatedWithSection = false;
|
||||
|
||||
DSAStack->checkMappableExprComponentListsForDeclAtLevel(
|
||||
@@ -1740,10 +1749,13 @@ bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const {
|
||||
|
||||
if (IsByRef && Ty.getNonReferenceType()->isScalarType()) {
|
||||
IsByRef =
|
||||
!DSAStack->hasExplicitDSA(
|
||||
D,
|
||||
[](OpenMPClauseKind K) -> bool { return K == OMPC_firstprivate; },
|
||||
Level, /*NotLastprivate=*/true) &&
|
||||
((IsVariableUsedInMapClause &&
|
||||
DSAStack->getCaptureRegion(Level, OpenMPCaptureLevel) ==
|
||||
OMPD_target) ||
|
||||
!DSAStack->hasExplicitDSA(
|
||||
D,
|
||||
[](OpenMPClauseKind K) -> bool { return K == OMPC_firstprivate; },
|
||||
Level, /*NotLastprivate=*/true)) &&
|
||||
// If the variable is artificial and must be captured by value - try to
|
||||
// capture by value.
|
||||
!(isa<OMPCapturedExprDecl>(D) && !D->hasAttr<OMPCaptureNoInitAttr>() &&
|
||||
@@ -2965,7 +2977,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
std::make_pair(StringRef(), QualType()) // __context with shared vars
|
||||
};
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
Params);
|
||||
Params, /*OpenMPCaptureLevel=*/0);
|
||||
// Mark this captured region as inlined, because we don't use outlined
|
||||
// function directly.
|
||||
getCurCapturedRegion()->TheCapturedDecl->addAttr(
|
||||
@@ -2976,7 +2988,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
};
|
||||
// Start a captured region for 'target' with no implicit parameters.
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
ParamsTarget);
|
||||
ParamsTarget, /*OpenMPCaptureLevel=*/1);
|
||||
Sema::CapturedParamNameType ParamsTeamsOrParallel[] = {
|
||||
std::make_pair(".global_tid.", KmpInt32PtrTy),
|
||||
std::make_pair(".bound_tid.", KmpInt32PtrTy),
|
||||
@@ -2985,7 +2997,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
// Start a captured region for 'teams' or 'parallel'. Both regions have
|
||||
// the same implicit parameters.
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
ParamsTeamsOrParallel);
|
||||
ParamsTeamsOrParallel, /*OpenMPCaptureLevel=*/2);
|
||||
break;
|
||||
}
|
||||
case OMPD_target:
|
||||
@@ -3009,14 +3021,15 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
std::make_pair(StringRef(), QualType()) // __context with shared vars
|
||||
};
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
Params);
|
||||
Params, /*OpenMPCaptureLevel=*/0);
|
||||
// Mark this captured region as inlined, because we don't use outlined
|
||||
// function directly.
|
||||
getCurCapturedRegion()->TheCapturedDecl->addAttr(
|
||||
AlwaysInlineAttr::CreateImplicit(
|
||||
Context, AlwaysInlineAttr::Keyword_forceinline));
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
std::make_pair(StringRef(), QualType()));
|
||||
std::make_pair(StringRef(), QualType()),
|
||||
/*OpenMPCaptureLevel=*/1);
|
||||
break;
|
||||
}
|
||||
case OMPD_simd:
|
||||
@@ -3148,7 +3161,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
std::make_pair(StringRef(), QualType()) // __context with shared vars
|
||||
};
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
Params);
|
||||
Params, /*OpenMPCaptureLevel=*/0);
|
||||
// Mark this captured region as inlined, because we don't use outlined
|
||||
// function directly.
|
||||
getCurCapturedRegion()->TheCapturedDecl->addAttr(
|
||||
@@ -3159,7 +3172,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
};
|
||||
// Start a captured region for 'target' with no implicit parameters.
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
ParamsTarget);
|
||||
ParamsTarget, /*OpenMPCaptureLevel=*/1);
|
||||
|
||||
Sema::CapturedParamNameType ParamsTeams[] = {
|
||||
std::make_pair(".global_tid.", KmpInt32PtrTy),
|
||||
@@ -3168,7 +3181,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
};
|
||||
// Start a captured region for 'target' with no implicit parameters.
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
ParamsTeams);
|
||||
ParamsTeams, /*OpenMPCaptureLevel=*/2);
|
||||
|
||||
Sema::CapturedParamNameType ParamsParallel[] = {
|
||||
std::make_pair(".global_tid.", KmpInt32PtrTy),
|
||||
@@ -3180,7 +3193,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
// Start a captured region for 'teams' or 'parallel'. Both regions have
|
||||
// the same implicit parameters.
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
ParamsParallel);
|
||||
ParamsParallel, /*OpenMPCaptureLevel=*/3);
|
||||
break;
|
||||
}
|
||||
|
||||
@@ -3197,7 +3210,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
};
|
||||
// Start a captured region for 'target' with no implicit parameters.
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
ParamsTeams);
|
||||
ParamsTeams, /*OpenMPCaptureLevel=*/0);
|
||||
|
||||
Sema::CapturedParamNameType ParamsParallel[] = {
|
||||
std::make_pair(".global_tid.", KmpInt32PtrTy),
|
||||
@@ -3209,7 +3222,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
|
||||
// Start a captured region for 'teams' or 'parallel'. Both regions have
|
||||
// the same implicit parameters.
|
||||
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
|
||||
ParamsParallel);
|
||||
ParamsParallel, /*OpenMPCaptureLevel=*/1);
|
||||
break;
|
||||
}
|
||||
case OMPD_target_update:
|
||||
@@ -11257,7 +11270,13 @@ OMPClause *Sema::ActOnOpenMPPrivateClause(ArrayRef<Expr *> VarList,
|
||||
// OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
|
||||
// A list item cannot appear in both a map clause and a data-sharing
|
||||
// attribute clause on the same construct
|
||||
if (isOpenMPTargetExecutionDirective(CurrDir)) {
|
||||
//
|
||||
// OpenMP 5.0 [2.19.7.1, Restrictions, p.7]
|
||||
// A list item cannot appear in both a map clause and a data-sharing
|
||||
// attribute clause on the same construct unless the construct is a
|
||||
// combined construct.
|
||||
if ((LangOpts.OpenMP <= 45 && isOpenMPTargetExecutionDirective(CurrDir)) ||
|
||||
CurrDir == OMPD_target) {
|
||||
OpenMPClauseKind ConflictKind;
|
||||
if (DSAStack->checkMappableExprComponentListsForDecl(
|
||||
VD, /*CurrentRegionOnly=*/true,
|
||||
@@ -11492,7 +11511,14 @@ OMPClause *Sema::ActOnOpenMPFirstprivateClause(ArrayRef<Expr *> VarList,
|
||||
// OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
|
||||
// A list item cannot appear in both a map clause and a data-sharing
|
||||
// attribute clause on the same construct
|
||||
if (isOpenMPTargetExecutionDirective(CurrDir)) {
|
||||
//
|
||||
// OpenMP 5.0 [2.19.7.1, Restrictions, p.7]
|
||||
// A list item cannot appear in both a map clause and a data-sharing
|
||||
// attribute clause on the same construct unless the construct is a
|
||||
// combined construct.
|
||||
if ((LangOpts.OpenMP <= 45 &&
|
||||
isOpenMPTargetExecutionDirective(CurrDir)) ||
|
||||
CurrDir == OMPD_target) {
|
||||
OpenMPClauseKind ConflictKind;
|
||||
if (DSAStack->checkMappableExprComponentListsForDecl(
|
||||
VD, /*CurrentRegionOnly=*/true,
|
||||
@@ -14606,7 +14632,14 @@ static void checkMappableExpressionList(
|
||||
// OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
|
||||
// A list item cannot appear in both a map clause and a data-sharing
|
||||
// attribute clause on the same construct
|
||||
if (VD && isOpenMPTargetExecutionDirective(DKind)) {
|
||||
//
|
||||
// OpenMP 5.0 [2.19.7.1, Restrictions, p.7]
|
||||
// A list item cannot appear in both a map clause and a data-sharing
|
||||
// attribute clause on the same construct unless the construct is a
|
||||
// combined construct.
|
||||
if (VD && ((SemaRef.LangOpts.OpenMP <= 45 &&
|
||||
isOpenMPTargetExecutionDirective(DKind)) ||
|
||||
DKind == OMPD_target)) {
|
||||
DSAStackTy::DSAVarData DVar = DSAS->getTopDSA(VD, /*FromParent=*/false);
|
||||
if (isOpenMPPrivate(DVar.CKind)) {
|
||||
SemaRef.Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
|
||||
|
||||
@@ -4334,7 +4334,8 @@ void Sema::ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
|
||||
|
||||
void Sema::ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
|
||||
CapturedRegionKind Kind,
|
||||
ArrayRef<CapturedParamNameType> Params) {
|
||||
ArrayRef<CapturedParamNameType> Params,
|
||||
unsigned OpenMPCaptureLevel) {
|
||||
CapturedDecl *CD = nullptr;
|
||||
RecordDecl *RD = CreateCapturedStmtRecordDecl(CD, Loc, Params.size());
|
||||
|
||||
@@ -4379,7 +4380,7 @@ void Sema::ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
|
||||
CD->setContextParam(ParamNum, Param);
|
||||
}
|
||||
// Enter the capturing scope for this captured region.
|
||||
PushCapturedRegionScope(CurScope, CD, RD, Kind);
|
||||
PushCapturedRegionScope(CurScope, CD, RD, Kind, OpenMPCaptureLevel);
|
||||
|
||||
if (CurScope)
|
||||
PushDeclContext(CurScope, CD);
|
||||
|
||||
@@ -1,6 +1,9 @@
|
||||
// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
|
||||
|
||||
// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
|
||||
|
||||
typedef void **omp_allocator_handle_t;
|
||||
extern const omp_allocator_handle_t omp_default_mem_alloc;
|
||||
@@ -240,7 +243,7 @@ int main(int argc, char **argv) {
|
||||
#pragma omp target teams distribute parallel for lastprivate(si) // OK
|
||||
for (i = 0; i < argc; ++i) si = i + 1;
|
||||
|
||||
#pragma omp target teams distribute parallel for lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for' directive}} expected-note {{defined as lastprivate}}
|
||||
#pragma omp target teams distribute parallel for lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for' directive}} le45-note {{defined as lastprivate}}
|
||||
for (i = 0; i < argc; ++i) foo();
|
||||
|
||||
return foomain<S4, S5>(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<S4, S5>' requested here}}
|
||||
|
||||
@@ -1,6 +1,9 @@
|
||||
// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
|
||||
|
||||
// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
|
||||
|
||||
typedef void **omp_allocator_handle_t;
|
||||
extern const omp_allocator_handle_t omp_default_mem_alloc;
|
||||
@@ -240,7 +243,7 @@ int main(int argc, char **argv) {
|
||||
#pragma omp target teams distribute parallel for simd lastprivate(si) // OK
|
||||
for (i = 0; i < argc; ++i) si = i + 1;
|
||||
|
||||
#pragma omp target teams distribute parallel for simd lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for simd' directive}} expected-note {{defined as lastprivate}}
|
||||
#pragma omp target teams distribute parallel for simd lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for simd' directive}} le45-note {{defined as lastprivate}}
|
||||
for (i = 0; i < argc; ++i) foo();
|
||||
return foomain<S4, S5>(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<S4, S5>' requested here}}
|
||||
}
|
||||
|
||||
@@ -1,6 +1,9 @@
|
||||
// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
|
||||
|
||||
// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
|
||||
|
||||
typedef void **omp_allocator_handle_t;
|
||||
extern const omp_allocator_handle_t omp_default_mem_alloc;
|
||||
@@ -131,8 +134,10 @@ int main(int argc, char **argv) {
|
||||
#pragma omp target teams distribute simd firstprivate(h) // expected-error {{threadprivate or thread local variable cannot be firstprivate}}
|
||||
for (i = 0; i < argc; ++i) foo();
|
||||
|
||||
#pragma omp target teams distribute simd private(i), firstprivate(i) // expected-error {{private variable cannot be firstprivate}} expected-note 2 {{defined as private}}
|
||||
for (i = 0; i < argc; ++i) foo(); // expected-error {{loop iteration variable in the associated loop of 'omp target teams distribute simd' directive may not be private, predetermined as linear}}
|
||||
#pragma omp target teams distribute simd private(i), firstprivate(i) // expected-error {{private variable cannot be firstprivate}} expected-note {{defined as private}}
|
||||
// le45-note@-1 {{defined as private}}
|
||||
// FIXME: Error is dropped for OpenMP 5.0. Probably shouldn't be.
|
||||
for (i = 0; i < argc; ++i) foo(); // le45-error {{loop iteration variable in the associated loop of 'omp target teams distribute simd' directive may not be private, predetermined as linear}}
|
||||
|
||||
#pragma omp target teams distribute simd firstprivate(i)
|
||||
for (j = 0; j < argc; ++j) foo();
|
||||
@@ -147,7 +152,7 @@ int main(int argc, char **argv) {
|
||||
#pragma omp target teams distribute simd lastprivate(argc), firstprivate(argc)
|
||||
for (i = 0; i < argc; ++i) foo();
|
||||
|
||||
#pragma omp target teams distribute simd firstprivate(argc) map(argc) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as firstprivate}}
|
||||
#pragma omp target teams distribute simd firstprivate(argc) map(argc) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as firstprivate}}
|
||||
for (i = 0; i < argc; ++i) foo();
|
||||
|
||||
return 0;
|
||||
|
||||
@@ -1,6 +1,9 @@
|
||||
// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
|
||||
|
||||
// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
|
||||
|
||||
typedef void **omp_allocator_handle_t;
|
||||
extern const omp_allocator_handle_t omp_default_mem_alloc;
|
||||
@@ -240,7 +243,7 @@ int main(int argc, char **argv) {
|
||||
#pragma omp target teams distribute simd lastprivate(si) // OK
|
||||
for (i = 0; i < argc; ++i) si = i + 1;
|
||||
|
||||
#pragma omp target teams distribute simd lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as lastprivate}}
|
||||
#pragma omp target teams distribute simd lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as lastprivate}}
|
||||
for (i = 0; i < argc; ++i) foo();
|
||||
|
||||
return foomain<S4, S5>(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<S4, S5>' requested here}}
|
||||
|
||||
@@ -1,6 +1,9 @@
|
||||
// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
|
||||
|
||||
// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
|
||||
|
||||
typedef void **omp_allocator_handle_t;
|
||||
extern const omp_allocator_handle_t omp_default_mem_alloc;
|
||||
@@ -132,7 +135,7 @@ int main(int argc, char **argv) {
|
||||
for (int k = 0; k < 10; ++k) {
|
||||
}
|
||||
|
||||
#pragma omp target teams distribute simd private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as private}}
|
||||
#pragma omp target teams distribute simd private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as private}}
|
||||
for (int k = 0; k < argc; ++k) ++k;
|
||||
|
||||
return 0;
|
||||
|
||||
172
clang/test/OpenMP/target_teams_map_codegen.cpp
Normal file
172
clang/test/OpenMP/target_teams_map_codegen.cpp
Normal file
@@ -0,0 +1,172 @@
|
||||
// Test host codegen.
|
||||
// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128
|
||||
// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST
|
||||
|
||||
// Test target codegen - host bc file has to be created first.
|
||||
// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
|
||||
// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefixes=CHECK,INT128
|
||||
// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
|
||||
// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,INT128
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
|
||||
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
|
||||
|
||||
// expected-no-diagnostics
|
||||
#ifndef HEADER
|
||||
#define HEADER
|
||||
|
||||
// HOST: @[[MAPTYPES_FIRSTPRIVATE:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
|
||||
// HOST: @[[MAPTYPES_REDUCTION:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
|
||||
// HOST: @[[MAPTYPES_FROM:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 34]
|
||||
// HOST: @[[MAPTYPES_TO:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 33]
|
||||
// HOST: @[[MAPTYPES_ALLOC:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 32]
|
||||
// HOST: @[[MAPTYPES_ARRAY_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
|
||||
// HOST: @[[MAPTYPES_ARRAY_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 33, i64 33]
|
||||
// HOST-INT128: @[[MAPTYPES_INT128_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
|
||||
// HOST-INT128: @[[MAPTYPES_INT128_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 34, i64 34]
|
||||
//
|
||||
// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_PRIVATE:__omp_offloading_[^"\\]*mapWithPrivate[^"\\]*]]\00"
|
||||
// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FIRSTPRIVATE:__omp_offloading_[^"\\]*mapWithFirstprivate[^"\\]*]]\00"
|
||||
// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_REDUCTION:__omp_offloading_[^"\\]*mapWithReduction[^"\\]*]]\00"
|
||||
// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FROM:__omp_offloading_[^"\\]*mapFrom[^"\\]*]]\00"
|
||||
// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_TO:__omp_offloading_[^"\\]*mapTo[^"\\]*]]\00"
|
||||
// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ALLOC:__omp_offloading_[^"\\]*mapAlloc[^"\\]*]]\00"
|
||||
// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R0:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00"
|
||||
// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R1:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00"
|
||||
// INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R0:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00"
|
||||
// INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R1:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00"
|
||||
|
||||
// HOST: define {{.*}}mapWithPrivate
|
||||
// HOST: call {{.*}} @.[[OFFLOAD_PRIVATE]].region_id
|
||||
// HOST-NOT: offload_maptypes
|
||||
// HOST-SAME: {{$}}
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OFFLOAD_PRIVATE]]()
|
||||
// CHECK: call void ({{.*}}@[[OUTLINE_PRIVATE:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OUTLINE_PRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid.)
|
||||
void mapWithPrivate() {
|
||||
int x, y;
|
||||
#pragma omp target teams private(x) map(x,y) private(y)
|
||||
;
|
||||
}
|
||||
|
||||
// HOST: define {{.*}}mapWithFirstprivate
|
||||
// HOST: call {{.*}} @.[[OFFLOAD_FIRSTPRIVATE]].region_id{{.*}} @[[MAPTYPES_FIRSTPRIVATE]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OFFLOAD_FIRSTPRIVATE]](i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y)
|
||||
// CHECK: call void ({{.*}}@[[OUTLINE_FIRSTPRIVATE:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OUTLINE_FIRSTPRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x, i{{[0-9]*}} %y)
|
||||
void mapWithFirstprivate() {
|
||||
int x, y;
|
||||
#pragma omp target teams firstprivate(x) map(x,y) firstprivate(y)
|
||||
;
|
||||
}
|
||||
|
||||
// HOST: define {{.*}}mapWithReduction
|
||||
// HOST: call {{.*}} @.[[OFFLOAD_REDUCTION]].region_id{{.*}} @[[MAPTYPES_REDUCTION]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OFFLOAD_REDUCTION]](i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y)
|
||||
// CHECK: call void ({{.*}}@[[OUTLINE_REDUCTION:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OUTLINE_REDUCTION]]({{[^)]*}} i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y)
|
||||
// CHECK: %.omp.reduction.red_list = alloca [2 x i8*]
|
||||
void mapWithReduction() {
|
||||
int x, y;
|
||||
#pragma omp target teams reduction(+:x) map(x,y) reduction(+:y)
|
||||
;
|
||||
}
|
||||
|
||||
// HOST: define {{.*}}mapFrom
|
||||
// HOST: call {{.*}} @.[[OFFLOAD_FROM]].region_id{{.*}} @[[MAPTYPES_FROM]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OFFLOAD_FROM]](i{{[0-9]*}}* {{[^,]*}}%x)
|
||||
// CHECK: call void ({{.*}}@[[OUTLINE_FROM:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OUTLINE_FROM]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x)
|
||||
void mapFrom() {
|
||||
int x;
|
||||
#pragma omp target teams firstprivate(x) map(from:x)
|
||||
;
|
||||
}
|
||||
|
||||
// HOST: define {{.*}}mapTo
|
||||
// HOST: call {{.*}} @.[[OFFLOAD_TO]].region_id{{.*}} @[[MAPTYPES_TO]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OFFLOAD_TO]](i{{[0-9]*}}* {{[^,]*}}%x)
|
||||
// CHECK: call void ({{.*}}@[[OUTLINE_TO:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OUTLINE_TO]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x)
|
||||
void mapTo() {
|
||||
int x;
|
||||
#pragma omp target teams firstprivate(x) map(to:x)
|
||||
;
|
||||
}
|
||||
|
||||
// HOST: define {{.*}}mapAlloc
|
||||
// HOST: call {{.*}} @.[[OFFLOAD_ALLOC]].region_id{{.*}} @[[MAPTYPES_ALLOC]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OFFLOAD_ALLOC]](i{{[0-9]*}}* {{[^,]*}}%x)
|
||||
// CHECK: call void ({{.*}}@[[OUTLINE_ALLOC:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OUTLINE_ALLOC]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x)
|
||||
void mapAlloc() {
|
||||
int x;
|
||||
#pragma omp target teams firstprivate(x) map(alloc:x)
|
||||
;
|
||||
}
|
||||
|
||||
// HOST: define {{.*}}mapArray
|
||||
// HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R0]].region_id{{.*}} @[[MAPTYPES_ARRAY_R0]]
|
||||
// HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R1]].region_id{{.*}} @[[MAPTYPES_ARRAY_R1]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R0]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
|
||||
// CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R0:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
|
||||
// CHECK: %.omp.reduction.red_list = alloca [1 x i8*]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R1]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
|
||||
// CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R1:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
|
||||
// CHECK: %.omp.reduction.red_list = alloca [1 x i8*]
|
||||
void mapArray() {
|
||||
int x[77], y[88], z[99];
|
||||
#pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z)
|
||||
;
|
||||
#pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(to:x,y,z)
|
||||
;
|
||||
}
|
||||
|
||||
# if HAS_INT128
|
||||
// HOST-INT128: define {{.*}}mapInt128
|
||||
// HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R0]].region_id{{.*}} @[[MAPTYPES_INT128_R0]]
|
||||
// HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R1]].region_id{{.*}} @[[MAPTYPES_INT128_R1]]
|
||||
//
|
||||
// INT128: define {{.*}} void @[[OFFLOAD_INT128_R0]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
|
||||
// INT128: call void ({{.*}}@[[OUTLINE_INT128_R0:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// INT128: define {{.*}} void @[[OUTLINE_INT128_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
|
||||
// INT128: %.omp.reduction.red_list = alloca [1 x i8*]
|
||||
//
|
||||
// INT128: define {{.*}} void @[[OFFLOAD_INT128_R1]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
|
||||
// INT128: call void ({{.*}}@[[OUTLINE_INT128_R1:.omp_outlined.[.0-9]*]]
|
||||
//
|
||||
// INT128: define {{.*}} void @[[OUTLINE_INT128_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
|
||||
// INT128: %.omp.reduction.red_list = alloca [1 x i8*]
|
||||
void mapInt128() {
|
||||
__int128 x, y, z;
|
||||
#pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z)
|
||||
;
|
||||
#pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(from:x,y,z)
|
||||
;
|
||||
}
|
||||
# endif
|
||||
#endif
|
||||
@@ -1,7 +1,10 @@
|
||||
// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
|
||||
|
||||
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
|
||||
// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-target -Wuninitialized
|
||||
// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
|
||||
// RUN: %clang_cc1 -DCCODE -verify=expected,le45 -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-target -Wuninitialized
|
||||
#ifdef CCODE
|
||||
void foo(int arg) {
|
||||
const int n = 0;
|
||||
@@ -536,10 +539,10 @@ int main(int argc, char **argv) {
|
||||
#pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
|
||||
foo();
|
||||
|
||||
#pragma omp target teams private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as private}}
|
||||
#pragma omp target teams private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}} le45-note {{defined as private}}
|
||||
{}
|
||||
|
||||
#pragma omp target teams firstprivate(j) map(j) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as firstprivate}}
|
||||
#pragma omp target teams firstprivate(j) map(j) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} le45-note {{defined as firstprivate}}
|
||||
{}
|
||||
|
||||
#pragma omp target teams map(m)
|
||||
|
||||
Reference in New Issue
Block a user