}
LValue ReductionCodeGen::emitSharedLValue(CodeGenFunction &CGF, const Expr *E) {
- if (const auto *OASE = dyn_cast<OMPArraySectionExpr>(E))
- return CGF.EmitOMPArraySectionExpr(OASE);
- if (const auto *ASE = dyn_cast<ArraySubscriptExpr>(E))
- return CGF.EmitLValue(ASE);
- auto *OrigVD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
- DeclRefExpr DRE(const_cast<VarDecl *>(OrigVD),
- CGF.CapturedStmtInfo &&
- CGF.CapturedStmtInfo->lookup(OrigVD) != nullptr,
- E->getType(), VK_LValue, E->getExprLoc());
- // Store the address of the original variable associated with the LHS
- // implicit variable.
- return CGF.EmitLValue(&DRE);
+ return CGF.EmitOMPSharedLValue(E);
}
LValue ReductionCodeGen::emitSharedLValueUB(CodeGenFunction &CGF,
OMP_MAP_PRIVATE_PTR = 0x80,
/// \brief Pass the element to the device by value.
OMP_MAP_PRIVATE_VAL = 0x100,
+ /// Implicit map
+ OMP_MAP_IMPLICIT = 0x200,
};
/// Class that associates information with a base pointer to be passed to the
OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
- bool IsFirstComponentList) const {
+ bool IsFirstComponentList, bool IsImplicit) const {
// The following summarizes what has to be generated for each map and the
// types bellow. The generated information is expressed in this order:
} else {
// The base is the reference to the variable.
// BP = &Var.
- BP = CGF.EmitLValue(cast<DeclRefExpr>(I->getAssociatedExpression()))
- .getPointer();
+ BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression()).getPointer();
// If the variable is a pointer and is being dereferenced (i.e. is not
// the last component), the base has to be the pointer itself, not its
}
}
+ unsigned DefaultFlags = IsImplicit ? OMP_MAP_IMPLICIT : 0;
for (; I != CE; ++I) {
auto Next = std::next(I);
isa<OMPArraySectionExpr>(Next->getAssociatedExpression())) &&
"Unexpected expression");
- auto *LB = CGF.EmitLValue(I->getAssociatedExpression()).getPointer();
+ llvm::Value *LB =
+ CGF.EmitOMPSharedLValue(I->getAssociatedExpression()).getPointer();
auto *Size = getExprTypeSize(I->getAssociatedExpression());
// If we have a member expression and the current component is a
BasePointers.push_back(BP);
Pointers.push_back(RefAddr);
Sizes.push_back(CGF.getTypeSize(CGF.getContext().VoidPtrTy));
- Types.push_back(getMapTypeBits(
- /*MapType*/ OMPC_MAP_alloc, /*MapTypeModifier=*/OMPC_MAP_unknown,
- !IsExpressionFirstInfo, IsCaptureFirstInfo));
+ Types.push_back(DefaultFlags |
+ getMapTypeBits(
+ /*MapType*/ OMPC_MAP_alloc,
+ /*MapTypeModifier=*/OMPC_MAP_unknown,
+ !IsExpressionFirstInfo, IsCaptureFirstInfo));
IsExpressionFirstInfo = false;
IsCaptureFirstInfo = false;
// The reference will be the next base address.
// same expression except for the first one. We also need to signal
// this map is the first one that relates with the current capture
// (there is a set of entries for each capture).
- Types.push_back(getMapTypeBits(MapType, MapTypeModifier,
- !IsExpressionFirstInfo,
- IsCaptureFirstInfo));
+ Types.push_back(DefaultFlags | getMapTypeBits(MapType, MapTypeModifier,
+ !IsExpressionFirstInfo,
+ IsCaptureFirstInfo));
// If we have a final array section, we are done with this expression.
if (IsFinalArraySection)
IsExpressionFirstInfo = false;
IsCaptureFirstInfo = false;
- continue;
}
}
}
RPK_MemberReference,
};
OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
- OpenMPMapClauseKind MapType;
- OpenMPMapClauseKind MapTypeModifier;
- ReturnPointerKind ReturnDevicePointer;
+ OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
+ OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown;
+ ReturnPointerKind ReturnDevicePointer = RPK_None;
+ bool IsImplicit = false;
- MapInfo()
- : MapType(OMPC_MAP_unknown), MapTypeModifier(OMPC_MAP_unknown),
- ReturnDevicePointer(RPK_None) {}
+ MapInfo() = default;
MapInfo(
OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapTypeModifier,
- ReturnPointerKind ReturnDevicePointer)
+ ReturnPointerKind ReturnDevicePointer, bool IsImplicit)
: Components(Components), MapType(MapType),
MapTypeModifier(MapTypeModifier),
- ReturnDevicePointer(ReturnDevicePointer) {}
+ ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit) {}
};
// We have to process the component lists that relate with the same
const ValueDecl *D,
OMPClauseMappableExprCommon::MappableExprComponentListRef L,
OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapModifier,
- MapInfo::ReturnPointerKind ReturnDevicePointer) {
+ MapInfo::ReturnPointerKind ReturnDevicePointer, bool IsImplicit) {
const ValueDecl *VD =
D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
- Info[VD].push_back({L, MapType, MapModifier, ReturnDevicePointer});
+ Info[VD].emplace_back(L, MapType, MapModifier, ReturnDevicePointer,
+ IsImplicit);
};
// FIXME: MSVC 2013 seems to require this-> to find member CurDir.
for (auto *C : this->CurDir.getClausesOfKind<OMPMapClause>())
- for (auto L : C->component_lists())
+ for (auto L : C->component_lists()) {
InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifier(),
- MapInfo::RPK_None);
+ MapInfo::RPK_None, C->isImplicit());
+ }
for (auto *C : this->CurDir.getClausesOfKind<OMPToClause>())
- for (auto L : C->component_lists())
+ for (auto L : C->component_lists()) {
InfoGen(L.first, L.second, OMPC_MAP_to, OMPC_MAP_unknown,
- MapInfo::RPK_None);
+ MapInfo::RPK_None, C->isImplicit());
+ }
for (auto *C : this->CurDir.getClausesOfKind<OMPFromClause>())
- for (auto L : C->component_lists())
+ for (auto L : C->component_lists()) {
InfoGen(L.first, L.second, OMPC_MAP_from, OMPC_MAP_unknown,
- MapInfo::RPK_None);
+ MapInfo::RPK_None, C->isImplicit());
+ }
// Look at the use_device_ptr clause information and mark the existing map
// entries as such. If there is no map information for an entry in the
// Remember the current base pointer index.
unsigned CurrentBasePointersIdx = BasePointers.size();
// FIXME: MSVC 2013 seems to require this-> to find the member method.
- this->generateInfoForComponentList(L.MapType, L.MapTypeModifier,
- L.Components, BasePointers, Pointers,
- Sizes, Types, IsFirstComponentList);
+ this->generateInfoForComponentList(
+ L.MapType, L.MapTypeModifier, L.Components, BasePointers, Pointers,
+ Sizes, Types, IsFirstComponentList, L.IsImplicit);
// If this entry relates with a device pointer, set the relevant
// declaration and add the 'return pointer' flag.
for (auto L : It->second) {
generateInfoForComponentList(
/*MapType=*/OMPC_MAP_to, /*MapTypeModifier=*/OMPC_MAP_unknown, L,
- BasePointers, Pointers, Sizes, Types, IsFirstComponentList);
+ BasePointers, Pointers, Sizes, Types, IsFirstComponentList,
+ /*IsImplicit=*/false);
IsFirstComponentList = false;
}
return;
"We got information for the wrong declaration??");
assert(!L.second.empty() &&
"Not expecting declaration with no component lists.");
- generateInfoForComponentList(C->getMapType(), C->getMapTypeModifier(),
- L.second, BasePointers, Pointers, Sizes,
- Types, IsFirstComponentList);
+ generateInfoForComponentList(
+ C->getMapType(), C->getMapTypeModifier(), L.second, BasePointers,
+ Pointers, Sizes, Types, IsFirstComponentList, C->isImplicit());
IsFirstComponentList = false;
}
// Stack of data-sharing attributes for variables
//===----------------------------------------------------------------------===//
+static Expr *CheckMapClauseExpressionBase(
+ Sema &SemaRef, Expr *E,
+ OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
+ OpenMPClauseKind CKind);
+
namespace {
/// \brief Default data sharing attributes, which can be applied to directive.
enum DefaultDataSharingAttributes {
bool ErrorFound;
CapturedStmt *CS;
llvm::SmallVector<Expr *, 8> ImplicitFirstprivate;
+ llvm::SmallVector<Expr *, 8> ImplicitMap;
llvm::DenseMap<ValueDecl *, Expr *> VarsWithInheritedDSA;
llvm::DenseSet<ValueDecl *> ImplicitDeclarations;
return;
}
+ if (isOpenMPTargetExecutionDirective(DKind) &&
+ !Stack->isLoopControlVariable(VD).first) {
+ if (!Stack->checkMappableExprComponentListsForDecl(
+ VD, /*CurrentRegionOnly=*/true,
+ [](OMPClauseMappableExprCommon::MappableExprComponentListRef
+ StackComponents,
+ OpenMPClauseKind) {
+ // Variable is used if it has been marked as an array, array
+ // section or the variable iself.
+ return StackComponents.size() == 1 ||
+ std::all_of(
+ std::next(StackComponents.rbegin()),
+ StackComponents.rend(),
+ [](const OMPClauseMappableExprCommon::
+ MappableComponent &MC) {
+ return MC.getAssociatedDeclaration() ==
+ nullptr &&
+ (isa<OMPArraySectionExpr>(
+ MC.getAssociatedExpression()) ||
+ isa<ArraySubscriptExpr>(
+ MC.getAssociatedExpression()));
+ });
+ })) {
+ bool CapturedByCopy = false;
+ // By default lambdas are captured as firstprivates.
+ if (const auto *RD =
+ VD->getType().getNonReferenceType()->getAsCXXRecordDecl())
+ if (RD->isLambda())
+ CapturedByCopy = true;
+ CapturedByCopy =
+ CapturedByCopy ||
+ llvm::any_of(
+ CS->captures(), [VD](const CapturedStmt::Capture &I) {
+ return I.capturesVariableByCopy() &&
+ I.getCapturedVar()->getCanonicalDecl() == VD;
+ });
+ if (CapturedByCopy)
+ ImplicitFirstprivate.emplace_back(E);
+ else
+ ImplicitMap.emplace_back(E);
+ return;
+ }
+ }
+
// OpenMP [2.9.3.6, Restrictions, p.2]
// A list item that appears in a reduction clause of the innermost
// enclosing worksharing or parallel construct may not be accessed in an
if (E->isTypeDependent() || E->isValueDependent() ||
E->containsUnexpandedParameterPack() || E->isInstantiationDependent())
return;
+ auto *FD = dyn_cast<FieldDecl>(E->getMemberDecl());
+ if (!FD)
+ return;
+ OpenMPDirectiveKind DKind = Stack->getCurrentDirective();
if (isa<CXXThisExpr>(E->getBase()->IgnoreParens())) {
- if (auto *FD = dyn_cast<FieldDecl>(E->getMemberDecl())) {
- auto DVar = Stack->getTopDSA(FD, false);
- // Check if the variable has explicit DSA set and stop analysis if it
- // so.
- if (DVar.RefExpr || !ImplicitDeclarations.insert(FD).second)
- return;
+ auto DVar = Stack->getTopDSA(FD, false);
+ // Check if the variable has explicit DSA set and stop analysis if it
+ // so.
+ if (DVar.RefExpr || !ImplicitDeclarations.insert(FD).second)
+ return;
- auto ELoc = E->getExprLoc();
- auto DKind = Stack->getCurrentDirective();
- // OpenMP [2.9.3.6, Restrictions, p.2]
- // A list item that appears in a reduction clause of the innermost
- // enclosing worksharing or parallel construct may not be accessed in
- // an explicit task.
- DVar = Stack->hasInnermostDSA(
- FD, [](OpenMPClauseKind C) -> bool { return C == OMPC_reduction; },
- [](OpenMPDirectiveKind K) -> bool {
- return isOpenMPParallelDirective(K) ||
- isOpenMPWorksharingDirective(K) ||
- isOpenMPTeamsDirective(K);
- },
- /*FromParent=*/true);
- if (isOpenMPTaskingDirective(DKind) && DVar.CKind == OMPC_reduction) {
- ErrorFound = true;
- SemaRef.Diag(ELoc, diag::err_omp_reduction_in_task);
- ReportOriginalDSA(SemaRef, Stack, FD, DVar);
+ if (isOpenMPTargetExecutionDirective(DKind) &&
+ !Stack->isLoopControlVariable(FD).first &&
+ !Stack->checkMappableExprComponentListsForDecl(
+ FD, /*CurrentRegionOnly=*/true,
+ [](OMPClauseMappableExprCommon::MappableExprComponentListRef
+ StackComponents,
+ OpenMPClauseKind) {
+ return isa<CXXThisExpr>(
+ cast<MemberExpr>(
+ StackComponents.back().getAssociatedExpression())
+ ->getBase()
+ ->IgnoreParens());
+ })) {
+ // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.3]
+ // A bit-field cannot appear in a map clause.
+ //
+ if (FD->isBitField()) {
+ SemaRef.Diag(E->getMemberLoc(),
+ diag::err_omp_bit_fields_forbidden_in_clause)
+ << E->getSourceRange() << getOpenMPClauseName(OMPC_map);
return;
}
+ ImplicitMap.emplace_back(E);
+ return;
+ }
- // Define implicit data-sharing attributes for task.
- DVar = Stack->getImplicitDSA(FD, false);
- if (isOpenMPTaskingDirective(DKind) && DVar.CKind != OMPC_shared &&
- !Stack->isLoopControlVariable(FD).first)
- ImplicitFirstprivate.push_back(E);
+ auto ELoc = E->getExprLoc();
+ // OpenMP [2.9.3.6, Restrictions, p.2]
+ // A list item that appears in a reduction clause of the innermost
+ // enclosing worksharing or parallel construct may not be accessed in
+ // an explicit task.
+ DVar = Stack->hasInnermostDSA(
+ FD, [](OpenMPClauseKind C) -> bool { return C == OMPC_reduction; },
+ [](OpenMPDirectiveKind K) -> bool {
+ return isOpenMPParallelDirective(K) ||
+ isOpenMPWorksharingDirective(K) || isOpenMPTeamsDirective(K);
+ },
+ /*FromParent=*/true);
+ if (isOpenMPTaskingDirective(DKind) && DVar.CKind == OMPC_reduction) {
+ ErrorFound = true;
+ SemaRef.Diag(ELoc, diag::err_omp_reduction_in_task);
+ ReportOriginalDSA(SemaRef, Stack, FD, DVar);
+ return;
+ }
+
+ // Define implicit data-sharing attributes for task.
+ DVar = Stack->getImplicitDSA(FD, false);
+ if (isOpenMPTaskingDirective(DKind) && DVar.CKind != OMPC_shared &&
+ !Stack->isLoopControlVariable(FD).first)
+ ImplicitFirstprivate.push_back(E);
+ return;
+ }
+ if (isOpenMPTargetExecutionDirective(DKind) && !FD->isBitField()) {
+ OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
+ CheckMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map);
+ auto *VD = cast<ValueDecl>(
+ CurComponents.back().getAssociatedDeclaration()->getCanonicalDecl());
+ if (!Stack->checkMappableExprComponentListsForDecl(
+ VD, /*CurrentRegionOnly=*/true,
+ [&CurComponents](
+ OMPClauseMappableExprCommon::MappableExprComponentListRef
+ StackComponents,
+ OpenMPClauseKind) {
+ if (CurComponents.size() < StackComponents.size())
+ return false;
+ auto CCI = CurComponents.rbegin();
+ for (const auto &SC : llvm::reverse(StackComponents)) {
+ // Do both expressions have the same kind?
+ if (CCI->getAssociatedExpression()->getStmtClass() !=
+ SC.getAssociatedExpression()->getStmtClass())
+ if (!(isa<OMPArraySectionExpr>(
+ SC.getAssociatedExpression()) &&
+ isa<ArraySubscriptExpr>(
+ CCI->getAssociatedExpression())))
+ return false;
+
+ Decl *CCD = CCI->getAssociatedDeclaration();
+ Decl *SCD = SC.getAssociatedDeclaration();
+ CCD = CCD ? CCD->getCanonicalDecl() : nullptr;
+ SCD = SCD ? SCD->getCanonicalDecl() : nullptr;
+ if (SCD != CCD)
+ return false;
+ std::advance(CCI, 1);
+ }
+ return true;
+ })) {
+ Visit(E->getBase());
}
} else
Visit(E->getBase());
void VisitOMPExecutableDirective(OMPExecutableDirective *S) {
for (auto *C : S->clauses()) {
// Skip analysis of arguments of implicitly defined firstprivate clause
- // for task directives.
- if (C && (!isa<OMPFirstprivateClause>(C) || C->getLocStart().isValid()))
+ // for task|target directives.
+ // Skip analysis of arguments of implicitly defined map clause for target
+ // directives.
+ if (C && !((isa<OMPFirstprivateClause>(C) || isa<OMPMapClause>(C)) &&
+ C->isImplicit())) {
for (auto *CC : C->children()) {
if (CC)
Visit(CC);
}
+ }
}
}
void VisitStmt(Stmt *S) {
}
bool isErrorFound() { return ErrorFound; }
- ArrayRef<Expr *> getImplicitFirstprivate() { return ImplicitFirstprivate; }
+ ArrayRef<Expr *> getImplicitFirstprivate() const {
+ return ImplicitFirstprivate;
+ }
+ ArrayRef<Expr *> getImplicitMap() const { return ImplicitMap; }
llvm::DenseMap<ValueDecl *, Expr *> &getVarsWithInheritedDSA() {
return VarsWithInheritedDSA;
}
SmallVector<Expr *, 4> ImplicitFirstprivates(
DSAChecker.getImplicitFirstprivate().begin(),
DSAChecker.getImplicitFirstprivate().end());
+ SmallVector<Expr *, 4> ImplicitMaps(DSAChecker.getImplicitMap().begin(),
+ DSAChecker.getImplicitMap().end());
// Mark taskgroup task_reduction descriptors as implicitly firstprivate.
for (auto *C : Clauses) {
if (auto *IRC = dyn_cast<OMPInReductionClause>(C)) {
} else
ErrorFound = true;
}
+ if (!ImplicitMaps.empty()) {
+ if (OMPClause *Implicit = ActOnOpenMPMapClause(
+ OMPC_MAP_unknown, OMPC_MAP_tofrom, /*IsMapTypeImplicit=*/true,
+ SourceLocation(), SourceLocation(), ImplicitMaps,
+ SourceLocation(), SourceLocation(), SourceLocation())) {
+ ClausesWithImplicit.emplace_back(Implicit);
+ ErrorFound |=
+ cast<OMPMapClause>(Implicit)->varlist_size() != ImplicitMaps.size();
+ } else
+ ErrorFound = true;
+ }
}
llvm::SmallVector<OpenMPDirectiveKind, 4> AllowedNameModifiers;
// CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK7-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
-// Map types: OMP_MAP_TO | OMP_MAP_IS_FIRST = 33
-// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33]
+// Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_IS_FIRST | OMP_MAP_IMPLICIT = 547
+// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547]
// CK7-LABEL: implicit_maps_double
void implicit_maps_double (int a){
#ifdef CK9
// CK9-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16]
-// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
-// CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 35]
+// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547
+// CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547]
// CK9-LABEL: implicit_maps_array
void implicit_maps_array (int a){
#ifdef CK11
// CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16]
-// Map types: OMP_MAP_TO + OMP_MAP_IS_FIRST = 33
-// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33]
+// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547
+// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547]
// CK11-LABEL: implicit_maps_double_complex
void implicit_maps_double_complex (int a){
// CK12-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
// Map types: OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
// CK12-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
-// Map types: OMP_MAP_TO + OMP_MAP_IS_FIRST = 33
-// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33]
+// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547
+// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547]
// CK12-LABEL: implicit_maps_float_complex
void implicit_maps_float_complex (int a){
// Map types:
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 (vla size)
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 (vla size)
-// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
-// CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 288, i32 288, i32 35]
+// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_IMPICIT_MAP = 547
+// CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 288, i32 288, i32 547]
// CK13-LABEL: implicit_maps_variable_length_array
void implicit_maps_variable_length_array (int a){
#ifdef CK14
// CK14-DAG: [[ST:%.+]] = type { i32, double }
-// CK14-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}, i{{64|32}} 4]
+// CK14-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i[[sz:64|32]]] [i{{64|32}} 4, i{{64|32}} 8, i{{64|32}} 4]
// Map types:
-// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
+// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_IMPLICIT_MAP = 547
+// - OMP_MAP_TO + OMP_MAP_FROM + OMP_IMPLICIT_MAP = 515
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
-// CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288]
+// CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 547, i32 515, i32 288]
class SSS {
public:
SSS sss(a, (double)a);
// CK14: define {{.*}}void @{{.+}}foo{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}})
- // CK14-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
+ // CK14-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 3, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
// CK14-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK14-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK14-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
// CK14-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
// CK14-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
- // CK14-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
+ // CK14-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
// CK14-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP0]]
- // CK14-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP0]]
+ // CK14-DAG: store i32* %{{.+}}, i32** [[CP0]]
// CK14-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
// CK14-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
+ // CK14-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+ // CK14-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
+ // CK14-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP1]]
+ // CK14-DAG: store double* %{{.+}}, double** [[CP1]]
+
+ // CK14-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2
+ // CK14-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2
// CK14-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz]]*
// CK14-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]*
// CK14-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP1]]
#ifdef CK15
// CK15: [[ST:%.+]] = type { i32, double, i32* }
-// CK15: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4]
+// CK15: [[SIZES:@.+]] = {{.+}}constant [5 x i[[sz:64|32]]] [i{{64|32}} 4, i{{64|32}} 8, i{{64|32}} {{8|4}}, i{{64|32}} 4, i{{64|32}} 4]
// Map types:
-// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
+// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
-// CK15: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288]
+// CK15: [[TYPES:@.+]] = {{.+}}constant [5 x i32] [i32 547, i32 515, i32 512, i32 531, i32 288]
-// CK15: [[SIZES2:@.+]] = {{.+}}constant [2 x i[[sz]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4]
+// CK15: [[SIZES2:@.+]] = {{.+}}constant [5 x i[[sz]]] [i{{64|32}} 4, i{{64|32}} 8, i{{64|32}} {{8|4}}, i{{64|32}} 4, i{{64|32}} 4]
// Map types:
// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
// - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
-// CK15: [[TYPES2:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288]
+// CK15: [[TYPES2:@.+]] = {{.+}}constant [5 x i32] [i32 547, i32 515, i32 512, i32 531, i32 288]
template<int x>
class SSST {
SSST<123> ssst(a, (double)a, a);
// CK15: define {{.*}}void @{{.+}}foo{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}})
- // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
+ // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 5, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
// CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
// CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
// CK15-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
- // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
+ // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
// CK15-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP0]]
- // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP0]]
+ // CK15-DAG: store i32* %{{.+}}, i32** [[CP0]]
// CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
// CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
- // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz]]*
- // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]*
- // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP1]]
- // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]]
+ // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+ // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
+ // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP1]]
+ // CK15-DAG: store double* %{{.+}}, double** [[CP1]]
+
+ // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2
+ // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2
+ // CK15-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+ // CK15-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32***
+ // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP2]]
+ // CK15-DAG: store i32** %{{.+}}, i32*** [[CP2]]
+
+ // CK15-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 3
+ // CK15-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 3
+ // CK15-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32***
+ // CK15-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
+ // CK15-DAG: store i32** %{{.+}}, i32*** [[CBP3]]
+ // CK15-DAG: store i32* %{{.+}}, i32** [[CP3]]
+
+ // CK15-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 4
+ // CK15-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 4
+ // CK15-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to i[[sz]]*
+ // CK15-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i[[sz]]*
+ // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP4]]
+ // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP4]]
// CK15-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
// CK15-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
// CK15-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
ssst.foo(456);
// CK15: define {{.*}}void @{{.+}}bar{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}})
- // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES2]]{{.+}}, {{.+}}[[TYPES2]]{{.+}})
+ // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 5, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES2]]{{.+}}, {{.+}}[[TYPES2]]{{.+}})
// CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
// CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
// CK15-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
- // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
- // CK15-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP0]]
- // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP0]]
+ // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+ // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP0]]
+ // CK15-DAG: store i32* %{{.+}}, i32** [[CP0]]
// CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
// CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
- // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz]]*
- // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]*
- // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP1]]
- // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]]
+ // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+ // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double**
+ // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP1]]
+ // CK15-DAG: store double* %{{.+}}, double** [[CP1]]
+
+ // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2
+ // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2
+ // CK15-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+ // CK15-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32***
+ // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP2]]
+ // CK15-DAG: store i32** %{{.+}}, i32*** [[CP2]]
+
+ // CK15-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 3
+ // CK15-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 3
+ // CK15-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32***
+ // CK15-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
+ // CK15-DAG: store i32** %{{.+}}, i32*** [[CBP3]]
+ // CK15-DAG: store i32* %{{.+}}, i32** [[CP3]]
+
+ // CK15-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 4
+ // CK15-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 4
+ // CK15-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to i[[sz]]*
+ // CK15-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i[[sz]]*
+ // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP4]]
+ // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP4]]
// CK15-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
// CK15-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
// CK15-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
// CK17-DAG: [[ST:%.+]] = type { i32, double }
// CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}]
-// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35
-// CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 35]
+// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547
+// CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547]
class SSS {
public:
// CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
-// CK24: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}]
-// CK24: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
-
-// CK24: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK24: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
-
-// CK24: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20]
-// CK24: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
-
-// CK24: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}]
-// CK24: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
-
-// CK24: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
-// CK24: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
-
-// CK24: [[SIZE07:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE07:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
-
-// CK24: [[SIZE08:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE08:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
-
-// CK24: [[SIZE09:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE09:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
-
-// CK24: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 8]
-// CK24: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
-
-// CK24: [[SIZE11:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}]
-// CK24: [[MTYPE11:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19]
-
-// CK24: [[SIZE12:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4]
-// CK24: [[MTYPE12:@.+]] = private {{.*}}constant [4 x i32] [i32 35, i32 19, i32 19, i32 19]
-
// CK24: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
#pragma omp target map(s.a)
{ s.a++; }
-// Region 02
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SA]]**
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store [[SA]]* [[SEC0:%.+]], [[SA]]** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24: call void [[CALL02:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.s.s)
- { s.a++; }
-
-// Region 03
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24: call void [[CALL03:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.s.s.a)
- { s.a++; }
-
-// Region 04
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 3
-
-// CK24: call void [[CALL04:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.b[:5])
- { s.a++; }
-
-// Region 05
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE05]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SB]]***
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store [[SB]]** [[SEC0:%.+]], [[SB]]*** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
-
-// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SB]]***
-// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[SB]]**
-// CK24-DAG: store [[SB]]** [[SEC0]], [[SB]]*** [[CBP1]]
-// CK24-DAG: store [[SB]]* [[SEC1:%.+]], [[SB]]** [[CP1]]
-// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0
-// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]],
-// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
-
-// CK24: call void [[CALL05:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.p[:5])
- { s.a++; }
-
-// Region 06
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE06]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[10 x [[SA]]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
-// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
-// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24: call void [[CALL06:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.s.sa[3].a)
- { s.a++; }
-
-// Region 07
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE07]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE07]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SA]]***
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store [[SA]]** [[SEC0:%.+]], [[SA]]*** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 3
-// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
-// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SA]]***
-// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
-// CK24-DAG: store [[SA]]** [[SEC0]], [[SA]]*** [[CBP1]]
-// CK24-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]]
-// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]],
-// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC1111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
-// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
-// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24: call void [[CALL07:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.s.sp[3]->a)
- { s.a++; }
-
-// Region 08
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE08]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SB]]***
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store [[SB]]** [[SEC0:%.+]], [[SB]]*** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
-
-// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SB]]***
-// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
-// CK24-DAG: store [[SB]]** [[SEC0]], [[SB]]*** [[CBP1]]
-// CK24-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]]
-// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0
-// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]],
-// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
-
-// CK24: call void [[CALL08:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.p->a)
- { s.a++; }
-
-// Region 09
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE09]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SA]]***
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store [[SA]]** [[SEC0:%.+]], [[SA]]*** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 4
-// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SA]]***
-// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
-// CK24-DAG: store [[SA]]** [[SEC0]], [[SA]]*** [[CBP1]]
-// CK24-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]]
-// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0
-// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]],
-// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SB]]* [[SEC1111:[^,]+]], i{{.+}} 0, i{{.+}} 4
-// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24: call void [[CALL09:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.s.p->a)
- { s.a++; }
-
-// Region 10
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SA]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
-// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24: call void [[CALL10:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.s.s.b[:2])
- { s.a++; }
-
-// Region 11
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE11]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SA]]***
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store [[SA]]** [[SEC0:%.+]], [[SA]]*** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 4
-// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SA]]***
-// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
-// CK24-DAG: store [[SA]]** [[SEC0]], [[SA]]*** [[CBP1]]
-// CK24-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]]
-// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[10 x i32]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[SEC11]] = getelementptr {{.*}}[[SA]]* [[SEC111:%[^,]+]], i{{.+}} 0, i{{.+}} 2
-// CK24-DAG: [[SEC111]] = load [[SA]]*, [[SA]]** [[SEC1111:%[^,]+]],
-// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 4
-// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1
-
-// CK24: call void [[CALL11:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.s.p->b[:2])
- { s.a++; }
-
-// Region 12
-// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE12]]{{.+}})
-// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
-// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
-
-// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]**
-// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SB]]***
-// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]]
-// CK24-DAG: store [[SB]]** [[SEC0:%.+]], [[SB]]*** [[CP0]]
-// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
-
-// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SB]]***
-// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[SA]]***
-// CK24-DAG: store [[SB]]** [[SEC0:%.+]], [[SB]]*** [[CBP1]]
-// CK24-DAG: store [[SA]]** [[SEC1:%.+]], [[SA]]*** [[CP1]]
-// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 4
-// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]],
-// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
-
-// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
-// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
-// CK24-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[SA]]***
-// CK24-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to [[SA]]***
-// CK24-DAG: store [[SA]]** [[SEC1:%.+]], [[SA]]*** [[CBP2]]
-// CK24-DAG: store [[SA]]** [[SEC2:%.+]], [[SA]]*** [[CP2]]
-// CK24-DAG: [[SEC2]] = getelementptr {{.*}}[[SA]]* [[SEC22:%[^,]+]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[SEC22]] = load [[SA]]*, [[SA]]** [[SEC222:%[^,]+]],
-// CK24-DAG: [[SEC222]] = getelementptr {{.*}}[[SB]]* [[SEC2222:%[^,]+]], i{{.+}} 0, i{{.+}} 4
-// CK24-DAG: [[SEC2222]] = load [[SB]]*, [[SB]]** [[SEC22222:%[^,]+]],
-// CK24-DAG: [[SEC22222]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
-
-// CK24-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
-// CK24-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
-// CK24-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to [[SA]]***
-// CK24-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
-// CK24-DAG: store [[SA]]** [[SEC2]], [[SA]]*** [[CBP3]]
-// CK24-DAG: store i32* [[SEC3:%.+]], i32** [[CP3]]
-// CK24-DAG: [[SEC3]] = getelementptr {{.*}}[[SA]]* [[SEC33:%[^,]+]], i{{.+}} 0, i{{.+}} 0
-// CK24-DAG: [[SEC33]] = load [[SA]]*, [[SA]]** [[SEC333:%[^,]+]],
-// CK24-DAG: [[SEC333]] = getelementptr {{.*}}[[SA]]* [[SEC3333:%[^,]+]], i{{.+}} 0, i{{.+}} 1
-// CK24-DAG: [[SEC3333]] = load [[SA]]*, [[SA]]** [[SEC33333:%[^,]+]],
-// CK24-DAG: [[SEC33333]] = getelementptr {{.*}}[[SB]]* [[SEC333333:%[^,]+]], i{{.+}} 0, i{{.+}} 4
-// CK24-DAG: [[SEC333333]] = load [[SB]]*, [[SB]]** [[SEC3333333:%[^,]+]],
-// CK24-DAG: [[SEC3333333]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2
-
-// CK24: call void [[CALL12:@.+]]([[SC]]* {{[^,]+}})
-#pragma omp target map(s.p->p->p->a)
- { s.a++; }
-
//
// Same thing but starting from a pointer.
//
}
// CK24: define {{.+}}[[CALL01]]
-// CK24: define {{.+}}[[CALL02]]
-// CK24: define {{.+}}[[CALL03]]
-// CK24: define {{.+}}[[CALL04]]
-// CK24: define {{.+}}[[CALL05]]
-// CK24: define {{.+}}[[CALL06]]
-// CK24: define {{.+}}[[CALL07]]
-// CK24: define {{.+}}[[CALL08]]
-// CK24: define {{.+}}[[CALL09]]
-// CK24: define {{.+}}[[CALL10]]
-// CK24: define {{.+}}[[CALL11]]
-// CK24: define {{.+}}[[CALL12]]
// CK24: define {{.+}}[[CALL13]]
// CK24: define {{.+}}[[CALL14]]
// CK24: define {{.+}}[[CALL15]]