Index: lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenMPRuntime.cpp +++ lib/CodeGen/CGOpenMPRuntime.cpp @@ -3781,27 +3781,481 @@ DeviceID, FileID, ParentName, Line, OutlinedFn, OutlinedFnID); } -void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF, - const OMPExecutableDirective &D, - llvm::Value *OutlinedFn, - llvm::Value *OutlinedFnID, - const Expr *IfCond, const Expr *Device, - ArrayRef CapturedVars) { - if (!CGF.HaveInsertPoint()) - return; +namespace { +// \brief Utility to extract information from the map clauses associated with a +// given construct and provide a convenient interface to obtain the information +// and generate code for that information. +class OpenMPMapClauseHandler { +public: /// \brief Values for bit flags used to specify the mapping type for /// offloading. enum OpenMPOffloadMappingFlags { + /// \brief Only allocate memory on the device, + OMP_MAP_ALLOC = 0x00, /// \brief Allocate memory on the device and move data from host to device. OMP_MAP_TO = 0x01, /// \brief Allocate memory on the device and move data from device to host. OMP_MAP_FROM = 0x02, + /// \brief Always perform the requested mapping action on the element, even + /// if it was already mapped before. + OMP_MAP_ALWAYS = 0x04, + /// \brief Decrement the reference count associated with the element without + /// executing any other action. + OMP_MAP_RELEASE = 0x08, + /// \brief Delete the element from the device environment, ignoring the + /// current reference count associated with the element. + OMP_MAP_DELETE = 0x10, /// \brief The element passed to the device is a pointer. OMP_MAP_PTR = 0x20, + /// \brief Signal the element as extra, i.e. is not argument to the target + /// region kernel. + OMP_MAP_EXTRA = 0x40, /// \brief Pass the element to the device by value. OMP_MAP_BYCOPY = 0x80, }; + typedef SmallVector MapValuesArrayTy; + typedef SmallVector MapFlagsArrayTy; + +private: + /// \brief Directive from where the map clauses were extracted. + const OMPExecutableDirective &Directive; + + /// \brief Function the directive is being generated for. + CodeGenFunction &CGF; + + struct DeclarationMapInfoEntry { + /// \brief Array of components in the map expression. + typedef std::pair ComponentTy; + typedef llvm::SmallVector ComponentsTy; + ComponentsTy Components; + + // Map type and modifier associated with this expression. + OpenMPMapClauseKind MapType; + OpenMPMapClauseKind MapTypeModifier; + + /// \brief Build and initialize this map information record with information + /// retrieved from the provided map clause expression. + DeclarationMapInfoEntry(const Expr *MCE, OpenMPMapClauseKind MapType, + OpenMPMapClauseKind MapTypeModifier) + : MapType(MapType), MapTypeModifier(MapTypeModifier) { + assert(MCE && "Invalid expression??"); + while (true) { + MCE = MCE->IgnoreParenImpCasts(); + + if (auto *CurE = dyn_cast(MCE)) { + Components.push_back( + ComponentTy(CurE, cast(CurE->getDecl()))); + break; + } + + if (auto *CurE = dyn_cast(MCE)) { + auto *BaseE = CurE->getBase()->IgnoreParenImpCasts(); + + Components.push_back( + ComponentTy(CurE, cast(CurE->getMemberDecl()))); + if (isa(BaseE)) + break; + + MCE = BaseE; + continue; + } + + if (auto *CurE = dyn_cast(MCE)) { + Components.push_back(ComponentTy(CurE, nullptr)); + MCE = CurE->getBase()->IgnoreParenImpCasts(); + continue; + } + + if (auto *CurE = dyn_cast(MCE)) { + Components.push_back(ComponentTy(CurE, nullptr)); + MCE = CurE->getBase()->IgnoreParenImpCasts(); + continue; + } + + llvm_unreachable("Invalid map clause expression!"); + } + } + + /// \brief Return declaration associated with this map information. If it is + /// a field it means the base is 'this'. + const ValueDecl *getAssociatedDecl() const { + assert(!Components.empty() && + "No expressions to extract declaration from??"); + const ValueDecl *D = Components.back().second; + assert(D && "Declaration must exist!"); + return D; + } + }; + + /// \brief Map between a declaration and its associated map information + /// entries. If the map info relates to 'this' we map it to null. + typedef SmallVector DeclarationMapInfoEntriesTy; + llvm::DenseMap + DeclarationMapInfoMap; + + llvm::Value *getExprTypeSize(const Expr *E) const { + auto ExprTy = E->getType().getCanonicalType(); + + // Reference types are ignored for mapping purposes. + if (auto *RefTy = ExprTy->getAs()) + ExprTy = RefTy->getPointeeType().getCanonicalType(); + + // Given that an array section is considered a built-in type, we need to + // do the calculation based on the length of the section instead of relying + // on CGF.getTypeSize(E->getType()). + if (const auto *OAE = dyn_cast(E)) { + auto BaseTy = + OAE->getBase()->IgnoreParenImpCasts()->getType().getCanonicalType(); + // Reference types are ignored for mapping purposes. + if (auto *RefTy = BaseTy->getAs()) + BaseTy = RefTy->getPointeeType().getCanonicalType(); + + // If there is no length associated with the expression, that means we + // are using the whole length of the base. + if (!OAE->getLength()) + return CGF.getTypeSize(BaseTy); + + llvm::Value *ElemSize; + if (auto *PTy = BaseTy->getAs()) { + ElemSize = CGF.getTypeSize(PTy->getPointeeType().getCanonicalType()); + } else { + auto *ATy = cast(BaseTy.getTypePtr()); + assert(ATy && "Expecting array type if not a pointer type."); + ElemSize = CGF.getTypeSize(ATy->getElementType().getCanonicalType()); + } + + auto *LengthVal = CGF.EmitScalarExpr(OAE->getLength()); + LengthVal = + CGF.Builder.CreateIntCast(LengthVal, CGF.SizeTy, /*isSigned=*/false); + return CGF.Builder.CreateNUWMul(LengthVal, ElemSize); + } + return CGF.getTypeSize(ExprTy); + } + + /// \brief Generate the address of the lower bound of the section defined by + /// expression \a E. + llvm::Value *getLowerBoundOfElement(const Expr *E) const { + return CGF.EmitLValue(E).getPointer(); + } + + /// \brief Return the corresponding bits for a given map clause modifier. Add + /// a flag marking the map as a pointer if requested. Add a flag marking the + /// map as extra, meaning is not an argument of the kernel. + unsigned getMapTypeBits(const DeclarationMapInfoEntry *Entry, bool AddPtrFlag, + bool AddExtraFlag) const { + unsigned Bits = 0u; + switch (Entry->MapType) { + case OMPC_MAP_alloc: + Bits = OMP_MAP_ALLOC; + break; + case OMPC_MAP_to: + Bits = OMP_MAP_TO; + break; + case OMPC_MAP_from: + Bits = OMP_MAP_FROM; + break; + case OMPC_MAP_tofrom: + Bits = OMP_MAP_TO | OMP_MAP_FROM; + break; + case OMPC_MAP_delete: + Bits = OMP_MAP_DELETE; + break; + case OMPC_MAP_release: + Bits = OMP_MAP_RELEASE; + break; + default: + llvm_unreachable("Unexpected map type!"); + break; + } + if (AddPtrFlag) + Bits |= OMP_MAP_PTR; + if (AddExtraFlag) + Bits |= OMP_MAP_EXTRA; + if (Entry->MapTypeModifier == OMPC_MAP_always) + Bits |= OMP_MAP_ALWAYS; + return Bits; + } + + /// \brief Generate the base pointers, section pointers, sizes and map type + /// bits for a given set of expressions \a MIE associated with a declaration. + void generateInfoForEntries(const DeclarationMapInfoEntriesTy &MIE, + MapValuesArrayTy &BasePointers, + MapValuesArrayTy &Pointers, + MapValuesArrayTy &Sizes, + MapFlagsArrayTy &Types) const { + // The following summarizes what has to be generated for each map and the + // types bellow. The generated information is expressed in this order: + // base pointer, section pointer, size, flags + // (to add to the ones that come from the map type and modifier). + // + // double d; + // int i[100]; + // float *p; + // + // struct S1 { + // int i; + // float f[50]; + // } + // struct S2 { + // int i; + // float f[50]; + // S1 s; + // double *p; + // struct S2 *ps; + // } + // S2 s; + // S2 *ps; + // + // map(d) + // &d, &d, sizeof(double), noflags + // + // map(i) + // &i, &i, 100*sizeof(int), noflags + // + // map(i[1:23]) + // &i(=&i[0]), &i[1], 23*sizeof(int), noflags + // + // map(p) + // &p, &p, sizeof(float*), noflags + // + // map(p[1:24]) + // p, &p[1], 24*sizeof(float), noflags + // + // map(s) + // &s, &s, sizeof(S2), noflags + // + // map(s.i) + // &s, &(s.i), sizeof(int), noflags + // + // map(s.s.f) + // &s, &(s.i.f), 50*sizeof(int), noflags + // + // map(s.p) + // &s, &(s.p), sizeof(double*), noflags + // + // map(s.p[:22], s.a s.b) + // &s, &(s.p), sizeof(double*), noflags + // &(s.p), &(s.p[0]), 22*sizeof(double), ptr_flag + extra_flag + // + // map(s.ps) + // &s, &(s.ps), sizeof(S2*), noflags + // + // map(s.ps->s.i) + // &s, &(s.ps), sizeof(S2*), noflags + // &(s.ps), &(s.ps->s.i), sizeof(int), ptr_flag + extra_flag + // + // map(s.ps->ps) + // &s, &(s.ps), sizeof(S2*), noflags + // &(s.ps), &(s.ps->ps), sizeof(S2*), ptr_flag + extra_flag + // + // map(s.ps->ps->ps) + // &s, &(s.ps), sizeof(S2*), noflags + // &(s.ps), &(s.ps->ps), sizeof(S2*), ptr_flag + extra_flag + // &(s.ps->ps), &(s.ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag + // + // map(s.ps->ps->s.f[:22]) + // &s, &(s.ps), sizeof(S2*), noflags + // &(s.ps), &(s.ps->ps), sizeof(S2*), ptr_flag + extra_flag + // &(s.ps->ps), &(s.ps->ps->s.f[0]), 22*sizeof(float), ptr_flag + extra_flag + // + // map(ps) + // &ps, &ps, sizeof(S2*), noflags + // + // map(ps->i) + // ps, &(ps->i), sizeof(int), noflags + // + // map(ps->s.f) + // ps, &(ps->s.f[0]), 50*sizeof(float), noflags + // + // map(ps->p) + // ps, &(ps->p), sizeof(double*), noflags + // + // map(ps->p[:22]) + // ps, &(ps->p), sizeof(double*), noflags + // &(ps->p), &(ps->p[0]), 22*sizeof(double), ptr_flag + extra_flag + // + // map(ps->ps) + // ps, &(ps->ps), sizeof(S2*), noflags + // + // map(ps->ps->s.i) + // ps, &(ps->ps), sizeof(S2*), noflags + // &(ps->ps), &(ps->ps->s.i), sizeof(int), ptr_flag + extra_flag + // + // map(ps->ps->ps) + // ps, &(ps->ps), sizeof(S2*), noflags + // &(ps->ps), &(ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag + // + // map(ps->ps->ps->ps) + // ps, &(ps->ps), sizeof(S2*), noflags + // &(ps->ps), &(ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag + // &(ps->ps->ps), &(ps->ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag + // + // map(ps->ps->ps->s.f[:22]) + // ps, &(ps->ps), sizeof(S2*), noflags + // &(ps->ps), &(ps->ps->ps), sizeof(S2*), ptr_flag + extra_flag + // &(ps->ps->ps), &(ps->ps->ps->s.f[0]), 22*sizeof(float), ptr_flag + + // extra_flag + + bool IsEntriesFirstInfo = true; + + // For each expression in the map clauses... + for (auto *InfoForExpr : MIE) { + auto CI = InfoForExpr->Components.rbegin(); + auto CE = InfoForExpr->Components.rend(); + auto I = CI; + + bool IsExpressionFirstInfo = true; + llvm::Value *BP = nullptr; + + if (auto *ME = dyn_cast(I->first)) { + // The base is the 'this' pointer. The content of the pointer is going + // to be the base of the field being mapped. + BP = CGF.EmitScalarExpr(ME->getBase()); + } else { + // The base is the reference to the variable. + // BP = &Var. + BP = CGF.EmitLValue(cast(I->first)).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 his + // reference. + if (I->second->getType()->isAnyPointerType() && std::next(I) != CE) { + auto PtrAddr = + CGF.MakeNaturalAlignAddrLValue(BP, I->second->getType()); + BP = CGF.EmitLoadOfLValue(PtrAddr, SourceLocation()).getScalarVal(); + + // We do not need to generate individual map information for the + // pointer, it can be associated with the combined storage. + ++I; + } + } + + for (; I != CE; ++I) { + auto Next = std::next(I); + + // We need to generate the addresses and sizes if this is the last + // component, or if the component is a pointer. In this case, the + // pointer becomes the base address for the following components. + if (Next == CE || I->first->getType()->isAnyPointerType()) { + + // If this is not the last component, we expect the pointer to be + // associated with an array expression or member expression. + assert((Next == CE || isa(Next->first) || + isa(Next->first) || + isa(Next->first)) && + "Unexpected expression"); + + // Save the base we are currently using. + BasePointers.push_back(BP); + + auto *LB = getLowerBoundOfElement(I->first); + auto *Size = getExprTypeSize(I->first); + + Pointers.push_back(LB); + Sizes.push_back(Size); + // We need to add a pointer flag for each map that comes from the the + // same expression except for the first one. We need to add the extra + // flag for each map that relates with the current capture, except for + // the first one (there is a set of entries for each capture). + Types.push_back(getMapTypeBits(InfoForExpr, !IsExpressionFirstInfo, + !IsEntriesFirstInfo)); + + // The pointer becomes the base for the next element. + if (Next != CE) + BP = LB; + + IsExpressionFirstInfo = false; + IsEntriesFirstInfo = false; + continue; + } + } + } + return; + } + +public: + OpenMPMapClauseHandler(const OMPExecutableDirective &Dir, + CodeGenFunction &CGF) + : Directive(Dir), CGF(CGF) { + + // Scan and extract information from the map clauses one by one. + for (auto *MC : Directive.getClausesOfKind()) { + for (auto *RE : MC->getVarRefs()) { + auto *MI = new DeclarationMapInfoEntry(RE, MC->getMapType(), + MC->getMapTypeModifier()); + const auto *D = MI->getAssociatedDecl(); + if (isa(D)) + DeclarationMapInfoMap[nullptr].push_back(MI); + else + DeclarationMapInfoMap[D].push_back(MI); + } + } + } + + ~OpenMPMapClauseHandler() { + // Clear the entries. + for (auto &MapEntries : DeclarationMapInfoMap) + for (auto *MapEntry : MapEntries.second) + delete (MapEntry); + } + + /// \brief Generate all the base pointers, section pointers, sizes and map + /// types for the extracted map information. + void generateAllInfo(MapValuesArrayTy &BasePointers, + MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, + MapFlagsArrayTy &Types) const { + BasePointers.clear(); + Pointers.clear(); + Sizes.clear(); + Types.clear(); + + // For each declaration identified in the map clause... + for (auto &MapEntries : DeclarationMapInfoMap) { + const DeclarationMapInfoEntriesTy &MIE = MapEntries.second; + generateInfoForEntries(MIE, BasePointers, Pointers, Sizes, Types); + } + return; + } + + /// \brief Generate the base pointers, section pointers, sizes and map types + /// associated to a given capture. + void generateInfoForCapture(const CapturedStmt::Capture *Cap, + MapValuesArrayTy &BasePointers, + MapValuesArrayTy &Pointers, + MapValuesArrayTy &Sizes, + MapFlagsArrayTy &Types) const { + assert(!Cap->capturesVariableArrayType() && + "Not expecting to generate map info for a variable array type!"); + + BasePointers.clear(); + Pointers.clear(); + Sizes.clear(); + Types.clear(); + + const ValueDecl *VD = Cap->capturesThis() ? nullptr : Cap->getCapturedVar(); + auto I = DeclarationMapInfoMap.find(VD); + if (I == DeclarationMapInfoMap.end()) + return; + + const DeclarationMapInfoEntriesTy &MIE = I->second; + assert(!MIE.empty() && "Not expecting declaration with empty information!"); + generateInfoForEntries(MIE, BasePointers, Pointers, Sizes, Types); + + return; + } +}; +} + +void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF, + const OMPExecutableDirective &D, + llvm::Value *OutlinedFn, + llvm::Value *OutlinedFnID, + const Expr *IfCond, const Expr *Device, + ArrayRef CapturedVars) { + if (!CGF.HaveInsertPoint()) + return; + enum OpenMPOffloadingReservedDeviceIDs { /// \brief Device ID if the device was not defined, runtime should get it /// from environment variables in the spec. @@ -3812,91 +4266,135 @@ auto &Ctx = CGF.getContext(); - // Fill up the arrays with the all the captured variables. - SmallVector BasePointers; - SmallVector Pointers; - SmallVector Sizes; - SmallVector MapTypes; + // Fill up the arrays with all the captured variables. + OpenMPMapClauseHandler::MapValuesArrayTy KernelArgs; + OpenMPMapClauseHandler::MapValuesArrayTy BasePointers; + OpenMPMapClauseHandler::MapValuesArrayTy Pointers; + OpenMPMapClauseHandler::MapValuesArrayTy Sizes; + OpenMPMapClauseHandler::MapFlagsArrayTy MapTypes; + + OpenMPMapClauseHandler::MapValuesArrayTy CurBasePointers; + OpenMPMapClauseHandler::MapValuesArrayTy CurPointers; + OpenMPMapClauseHandler::MapValuesArrayTy CurSizes; + OpenMPMapClauseHandler::MapFlagsArrayTy CurMapTypes; - bool hasVLACaptures = false; + // Get map clause information. + OpenMPMapClauseHandler MCHandler(D, CGF); const CapturedStmt &CS = *cast(D.getAssociatedStmt()); auto RI = CS.getCapturedRecordDecl()->field_begin(); - // auto II = CS.capture_init_begin(); auto CV = CapturedVars.begin(); for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(), CE = CS.capture_end(); CI != CE; ++CI, ++RI, ++CV) { StringRef Name; QualType Ty; - llvm::Value *BasePointer; - llvm::Value *Pointer; - llvm::Value *Size; - unsigned MapType; - // VLA sizes are passed to the outlined region by copy. + CurBasePointers.clear(); + CurPointers.clear(); + CurSizes.clear(); + CurMapTypes.clear(); + + // VLA sizes are passed to the outlined region by copy and do not have map + // information associated. if (CI->capturesVariableArrayType()) { - BasePointer = Pointer = *CV; - Size = CGF.getTypeSize(RI->getType()); + CurBasePointers.push_back(*CV); + CurPointers.push_back(*CV); + CurSizes.push_back(CGF.getTypeSize(RI->getType())); // Copy to the device as an argument. No need to retrieve it. - MapType = OMP_MAP_BYCOPY; - hasVLACaptures = true; - } else if (CI->capturesThis()) { - BasePointer = Pointer = *CV; - const PointerType *PtrTy = cast(RI->getType().getTypePtr()); - Size = CGF.getTypeSize(PtrTy->getPointeeType()); - // Default map type. - MapType = OMP_MAP_TO | OMP_MAP_FROM; - } else if (CI->capturesVariableByCopy()) { - MapType = OMP_MAP_BYCOPY; - if (!RI->getType()->isAnyPointerType()) { - // If the field is not a pointer, we need to save the actual value and - // load it as a void pointer. - auto DstAddr = CGF.CreateMemTemp( - Ctx.getUIntPtrType(), - Twine(CI->getCapturedVar()->getName()) + ".casted"); - LValue DstLV = CGF.MakeAddrLValue(DstAddr, Ctx.getUIntPtrType()); - - auto *SrcAddrVal = CGF.EmitScalarConversion( - DstAddr.getPointer(), Ctx.getPointerType(Ctx.getUIntPtrType()), - Ctx.getPointerType(RI->getType()), SourceLocation()); - LValue SrcLV = - CGF.MakeNaturalAlignAddrLValue(SrcAddrVal, RI->getType()); - - // Store the value using the source type pointer. - CGF.EmitStoreThroughLValue(RValue::get(*CV), SrcLV); - - // Load the value using the destination type pointer. - BasePointer = Pointer = - CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal(); - } else { - MapType |= OMP_MAP_PTR; - BasePointer = Pointer = *CV; - } - Size = CGF.getTypeSize(RI->getType()); + CurMapTypes.push_back(OpenMPMapClauseHandler::OMP_MAP_BYCOPY); } else { - assert(CI->capturesVariable() && "Expected captured reference."); - BasePointer = Pointer = *CV; - - const ReferenceType *PtrTy = - cast(RI->getType().getTypePtr()); - QualType ElementType = PtrTy->getPointeeType(); - Size = CGF.getTypeSize(ElementType); - // The default map type for a scalar/complex type is 'to' because by - // default the value doesn't have to be retrieved. For an aggregate type, - // the default is 'tofrom'. - MapType = ElementType->isAggregateType() ? (OMP_MAP_TO | OMP_MAP_FROM) - : OMP_MAP_TO; - if (ElementType->isAnyPointerType()) - MapType |= OMP_MAP_PTR; + // If we have any information in the map clause, we use it, otherwise we + // just do a default mapping. + MCHandler.generateInfoForCapture(CI, CurBasePointers, CurPointers, + CurSizes, CurMapTypes); + + if (CurBasePointers.empty()) { + // Do the default mapping. + if (CI->capturesThis()) { + CurBasePointers.push_back(*CV); + CurPointers.push_back(*CV); + const PointerType *PtrTy = + cast(RI->getType().getTypePtr()); + CurSizes.push_back(CGF.getTypeSize(PtrTy->getPointeeType())); + // Default map type. + CurMapTypes.push_back(OpenMPMapClauseHandler::OMP_MAP_TO | + OpenMPMapClauseHandler::OMP_MAP_FROM); + } else if (CI->capturesVariableByCopy()) { + CurMapTypes.push_back(OpenMPMapClauseHandler::OMP_MAP_BYCOPY); + if (!RI->getType()->isAnyPointerType()) { + // If the field is not a pointer, we need to save the actual value + // and + // load it as a void pointer. + auto DstAddr = CGF.CreateMemTemp( + Ctx.getUIntPtrType(), + Twine(CI->getCapturedVar()->getName()) + ".casted"); + LValue DstLV = CGF.MakeAddrLValue(DstAddr, Ctx.getUIntPtrType()); + + auto *SrcAddrVal = CGF.EmitScalarConversion( + DstAddr.getPointer(), Ctx.getPointerType(Ctx.getUIntPtrType()), + Ctx.getPointerType(RI->getType()), SourceLocation()); + LValue SrcLV = + CGF.MakeNaturalAlignAddrLValue(SrcAddrVal, RI->getType()); + + // Store the value using the source type pointer. + CGF.EmitStoreThroughLValue(RValue::get(*CV), SrcLV); + + // Load the value using the destination type pointer. + CurBasePointers.push_back( + CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal()); + CurPointers.push_back(CurBasePointers.back()); + } else { + CurBasePointers.push_back(*CV); + CurPointers.push_back(*CV); + } + CurSizes.push_back(CGF.getTypeSize(RI->getType())); + } else { + assert(CI->capturesVariable() && "Expected captured reference."); + CurBasePointers.push_back(*CV); + CurPointers.push_back(*CV); + + const ReferenceType *PtrTy = + cast(RI->getType().getTypePtr()); + QualType ElementType = PtrTy->getPointeeType(); + CurSizes.push_back(CGF.getTypeSize(ElementType)); + // The default map type for a scalar/complex type is 'to' because by + // default the value doesn't have to be retrieved. For an aggregate + // type, + // the default is 'tofrom'. + CurMapTypes.push_back(ElementType->isAggregateType() + ? (OpenMPMapClauseHandler::OMP_MAP_TO | + OpenMPMapClauseHandler::OMP_MAP_FROM) + : OpenMPMapClauseHandler::OMP_MAP_TO); + } + } } - - BasePointers.push_back(BasePointer); - Pointers.push_back(Pointer); - Sizes.push_back(Size); - MapTypes.push_back(MapType); + // We expect to have at least an element of information for this capture. + assert(!CurBasePointers.empty() && "Non-existing map pointer for capture!"); + assert(CurBasePointers.size() == CurPointers.size() && + CurBasePointers.size() == CurSizes.size() && + CurBasePointers.size() == CurMapTypes.size() && + "Inconsistent map information sizes!"); + + // The kernel args are always the first elements of the base pointers + // associated with a capture. + KernelArgs.push_back(CurBasePointers.front()); + // We need to append the results of this capture to what we already have. + BasePointers.append(CurBasePointers.begin(), CurBasePointers.end()); + Pointers.append(CurPointers.begin(), CurPointers.end()); + Sizes.append(CurSizes.begin(), CurSizes.end()); + MapTypes.append(CurMapTypes.begin(), CurMapTypes.end()); } + // Detect if we have any capture size requiring runtime evaluation of the size + // so that a constant array could be eventually used. + bool hasRuntimeEvaluationCaptureSize = false; + for (auto *S : Sizes) + if (!isa(S)) { + hasRuntimeEvaluationCaptureSize = true; + break; + } + // Keep track on whether the host function has to be executed. auto OffloadErrorQType = Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/true); @@ -3908,8 +4406,8 @@ // Fill up the pointer arrays and transfer execution to the device. auto &&ThenGen = [this, &Ctx, &BasePointers, &Pointers, &Sizes, &MapTypes, - hasVLACaptures, Device, OutlinedFnID, OffloadError, - OffloadErrorQType](CodeGenFunction &CGF) { + hasRuntimeEvaluationCaptureSize, Device, OutlinedFnID, + OffloadError, OffloadErrorQType](CodeGenFunction &CGF) { unsigned PointerNumVal = BasePointers.size(); llvm::Value *PointerNum = CGF.Builder.getInt32(PointerNumVal); llvm::Value *BasePointersArray; @@ -3931,7 +4429,7 @@ // If we don't have any VLA types, we can use a constant array for the map // sizes, otherwise we need to fill up the arrays as we do for the // pointers. - if (hasVLACaptures) { + if (hasRuntimeEvaluationCaptureSize) { QualType SizeArrayType = Ctx.getConstantArrayType( Ctx.getSizeType(), PointerNumAP, ArrayType::Normal, /*IndexTypeQuals=*/0); @@ -3994,7 +4492,7 @@ Address PAddr(P, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy)); CGF.Builder.CreateStore(PVal, PAddr); - if (hasVLACaptures) { + if (hasRuntimeEvaluationCaptureSize) { llvm::Value *S = CGF.Builder.CreateConstInBoundsGEP2_32( llvm::ArrayType::get(CGM.SizeTy, PointerNumVal), SizesArray, /*Idx0=*/0, @@ -4089,7 +4587,7 @@ CGF.Builder.CreateCondBr(Failed, OffloadFailedBlock, OffloadContBlock); CGF.EmitBlock(OffloadFailedBlock); - CGF.Builder.CreateCall(OutlinedFn, BasePointers); + CGF.Builder.CreateCall(OutlinedFn, KernelArgs); CGF.EmitBranch(OffloadContBlock); CGF.EmitBlock(OffloadContBlock, /*IsFinished=*/true); Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -809,6 +809,55 @@ #define DSAStack static_cast(VarDataSharingAttributesStack) +namespace { +// Types used to organize the components of a valid map clause. +typedef std::pair MapExpressionComponent; +typedef SmallVector MapExpressionComponents; +} + +// Helper to extract the components in the map clause expression \a E and store +// them into \a MEC. This assumes that \a E is a valid map clause expression, +// i.e. it has already passed the single clause checks. +static void ExtractMapExpressionComponents(Expr *TE, + MapExpressionComponents &MEC) { + while (true) { + TE = TE->IgnoreParenImpCasts(); + + if (auto *CurE = dyn_cast(TE)) { + MEC.push_back( + MapExpressionComponent(CurE, cast(CurE->getDecl()))); + break; + } + + if (auto *CurE = dyn_cast(TE)) { + auto *BaseE = CurE->getBase()->IgnoreParenImpCasts(); + + MEC.push_back( + MapExpressionComponent(CurE, cast(CurE->getMemberDecl()))); + if (isa(BaseE)) + break; + + TE = BaseE; + continue; + } + + if (auto *CurE = dyn_cast(TE)) { + MEC.push_back(MapExpressionComponent(CurE, nullptr)); + TE = CurE->getBase()->IgnoreParenImpCasts(); + continue; + } + + if (auto *CurE = dyn_cast(TE)) { + MEC.push_back(MapExpressionComponent(CurE, nullptr)); + TE = CurE->getBase()->IgnoreParenImpCasts(); + continue; + } + + llvm_unreachable( + "Expecting only valid map clause expressions at this point!"); + } +} + bool Sema::IsOpenMPCapturedByRef(ValueDecl *D, const CapturedRegionScopeInfo *RSI) { assert(LangOpts.OpenMP && "OpenMP is not allowed"); @@ -875,16 +924,60 @@ // array section, the runtime library may pass the NULL value to the // device instead of the value passed to it by the compiler. - // FIXME: Right now, only implicit maps are implemented. Properly mapping - // values requires having the map, private, and firstprivate clauses SEMA - // and parsing in place, which we don't yet. if (Ty->isReferenceType()) Ty = Ty->castAs()->getPointeeType(); - IsByRef = !Ty->isScalarType(); + + // FIXME: Right now, only some of the maps are implemented. Properly mapping + // values requires having the private, and firstprivate clauses SEMA + // and parsing in place, which we don't yet. + + // 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->checkMapInfoForVar( + D, /*CurrentRegionOnly=*/true, [&](Expr *MapExpr) { + MapExpressionComponents MapExprComponents; + ExtractMapExpressionComponents(MapExpr, MapExprComponents); + + auto EI = MapExprComponents.rbegin(); + auto EE = MapExprComponents.rend(); + + assert(EI != EE && "Invalid map expression!"); + + if (isa(EI->first)) + IsVariableUsedInMapClause |= EI->second == D; + + ++EI; + if (EI == EE) + return false; + + if (isa(EI->first) || + isa(EI->first) || + isa(EI->first)) { + IsVariableAssociatedWithSection = true; + // There is nothing more we need to know about this variable. + return true; + } + + // Keep looking for more map info. + return false; + }); + + if (IsVariableUsedInMapClause) { + // If variable is identified in a map clause it is always captured by + // reference except if it is a pointer that is dereferenced somehow. + IsByRef = !(Ty->isPointerType() && IsVariableAssociatedWithSection); + } else { + // By default, all the data that has a scalar type is mapped by copy. + IsByRef = !Ty->isScalarType(); + } } - // When passing data by value, we need to make sure it fits the uintptr size + // When passing data by copy, we need to make sure it fits the uintptr size // and alignment, because the runtime library only deals with uintptr types. // If it does not fit the uintptr size, we need to pass the data by reference // instead. @@ -9108,54 +9201,6 @@ static bool CheckMapConflicts(Sema &SemaRef, DSAStackTy *DSAS, ValueDecl *VD, Expr *E, bool CurrentRegionOnly) { assert(VD && E); - - // Types used to organize the components of a valid map clause. - typedef std::pair MapExpressionComponent; - typedef SmallVector MapExpressionComponents; - - // Helper to extract the components in the map clause expression E and store - // them into MEC. This assumes that E is a valid map clause expression, i.e. - // it has already passed the single clause checks. - auto ExtractMapExpressionComponents = [](Expr *TE, - MapExpressionComponents &MEC) { - while (true) { - TE = TE->IgnoreParenImpCasts(); - - if (auto *CurE = dyn_cast(TE)) { - MEC.push_back( - MapExpressionComponent(CurE, cast(CurE->getDecl()))); - break; - } - - if (auto *CurE = dyn_cast(TE)) { - auto *BaseE = CurE->getBase()->IgnoreParenImpCasts(); - - MEC.push_back(MapExpressionComponent( - CurE, cast(CurE->getMemberDecl()))); - if (isa(BaseE)) - break; - - TE = BaseE; - continue; - } - - if (auto *CurE = dyn_cast(TE)) { - MEC.push_back(MapExpressionComponent(CurE, nullptr)); - TE = CurE->getBase()->IgnoreParenImpCasts(); - continue; - } - - if (auto *CurE = dyn_cast(TE)) { - MEC.push_back(MapExpressionComponent(CurE, nullptr)); - TE = CurE->getBase()->IgnoreParenImpCasts(); - continue; - } - - llvm_unreachable( - "Expecting only valid map clause expressions at this point!"); - } - }; - SourceLocation ELoc = E->getExprLoc(); SourceRange ERange = E->getSourceRange(); Index: test/OpenMP/target_map_codegen.cpp =================================================================== --- test/OpenMP/target_map_codegen.cpp +++ test/OpenMP/target_map_codegen.cpp @@ -454,8 +454,8 @@ #ifdef CK10 // CK10-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}] -// Map types: OMP_MAP_BYCOPY | OMP_MAP_PTR = 128 + 32 -// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 160] +// Map types: OMP_MAP_BYCOPY = 128 +// CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] // CK10-LABEL: implicit_maps_pointer void implicit_maps_pointer (){ @@ -1012,4 +1012,2464 @@ // CK18-32: {{.+}} = load i32, i32* [[ADDR]], #endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK19 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64 +// RUN: %clang_cc1 -DCK19 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK19 --check-prefix CK19-64 +// RUN: %clang_cc1 -DCK19 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-32 +// RUN: %clang_cc1 -DCK19 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK19 --check-prefix CK19-32 +#ifdef CK19 + +// CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK19: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK19: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK19: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240] +// CK19: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK19: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240] +// CK19: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK19: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK19: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK19: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240] +// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 240] +// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] zeroinitializer + +// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK19: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 1] + +// CK19: [[SIZE17:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 240] +// CK19: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 2] + +// CK19: [[SIZE18:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 240] +// CK19: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3] + +// CK19: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 0] + +// CK19: [[SIZE20:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK19: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 1] + +// CK19: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3] + +// CK19: [[SIZE22:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK19: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i32] [i32 128, i32 3] + +// CK19: [[SIZE23:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i32] [i32 7] + +// CK19: [[SIZE24:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 480] +// CK19: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE25:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK19: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE26:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 24] +// CK19: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE27:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK19: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK19: [[SIZE28:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 16] +// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99] + +// CK19: [[SIZE29:@.+]] = private {{.*}}constant [3 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4] +// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i32] [i32 3, i32 99, i32 99] + +// CK19-LABEL: explicit_maps_single +void explicit_maps_single (int ii){ + + // Map of a scalar. + int a = ii; + + // Region 00 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8* + + // CK19: call void [[CALL00:@.+]](i32* {{[^,]+}}) + #pragma omp target map(alloc:a) + { + ++a; + } + + // Map of an array. + int arra[100]; + + // Region 01 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [100 x i32]* [[VAR0]] to i8* + + // CK19: call void [[CALL01:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(to:arra) + { + arra[50]++; + } + + // Region 02 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 20 + + // CK19: call void [[CALL02:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(from:arra[20:60]) + { + arra[50]++; + } + + // Region 03 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + + // CK19: call void [[CALL03:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(tofrom:arra[:60]) + { + arra[50]++; + } + + // Region 04 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + + // CK19: call void [[CALL04:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(alloc:arra[:]) + { + arra[50]++; + } + + // Region 05 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 15 + + // CK19: call void [[CALL05:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(to:arra[15]) + { + arra[15]++; + } + + // Region 06 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}} + + // CK19: call void [[CALL06:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(tofrom:arra[ii:ii+23]) + { + arra[50]++; + } + + // Region 07 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE07]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + + // CK19: call void [[CALL07:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(alloc:arra[:ii]) + { + arra[50]++; + } + + // Region 08 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE08]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [100 x i32]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}} + + // CK19: call void [[CALL08:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(tofrom:arra[ii]) + { + arra[15]++; + } + + // Map of a pointer. + int *pa; + + // Region 09 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32** [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32** [[VAR0]] to i8* + + // CK19: call void [[CALL09:@.+]](i32** {{[^,]+}}) + #pragma omp target map(from:pa) + { + pa[50]++; + } + + // Region 10 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 20 + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL10:@.+]](i32* {{[^,]+}}) + #pragma omp target map(tofrom:pa[20:60]) + { + pa[50]++; + } + + // Region 11 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE11]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL11:@.+]](i32* {{[^,]+}}) + #pragma omp target map(alloc:pa[:60]) + { + pa[50]++; + } + + // Region 12 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE12]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 15 + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL12:@.+]](i32* {{[^,]+}}) + #pragma omp target map(to:pa[15]) + { + pa[15]++; + } + + // Region 13 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}} + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL13:@.+]](i32* {{[^,]+}}) + #pragma omp target map(alloc:pa[ii-23:ii]) + { + pa[50]++; + } + + // Region 14 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[CSVAL0]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL14:@.+]](i32* {{[^,]+}}) + #pragma omp target map(to:pa[:ii]) + { + pa[50]++; + } + + // Region 15 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE15]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE15]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}} + // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK19: call void [[CALL15:@.+]](i32* {{[^,]+}}) + #pragma omp target map(from:pa[ii+12]) + { + pa[15]++; + } + + // Map of a variable-size array. + int va[ii]; + + // Region 16 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE16]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[VAR1]] to i8* + // CK19-DAG: [[CSVAL1]] = mul nuw i{{.+}} %{{.*}}, 4 + + // CK19: call void [[CALL16:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(to:va) + { + va[50]++; + } + + // Region 17 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE17]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE17]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 20 + + // CK19: call void [[CALL17:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(from:va[20:60]) + { + va[50]++; + } + + // Region 18 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE18]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE18]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 0 + + // CK19: call void [[CALL18:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(tofrom:va[:60]) + { + va[50]++; + } + + // Region 19 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE19]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[CSVAL1]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 0 + + // CK19: call void [[CALL19:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(alloc:va[:]) + { + va[50]++; + } + + // Region 20 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE20]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE20]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 15 + + // CK19: call void [[CALL20:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(to:va[15]) + { + va[15]++; + } + + // Region 21 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE21]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[S0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[CSVAL1]] = mul nuw i{{.+}} %{{.*}}, 4 + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} %{{.+}} + + // CK19: call void [[CALL21:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(tofrom:va[ii:ii+23]) + { + va[50]++; + } + + // Region 22 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE22]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE22]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = inttoptr i[[Z]] %{{.+}} to i8* + // CK19-DAG: [[CPVAL0]] = inttoptr i[[Z]] %{{.+}}to i8* + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32* [[VAR1:%.+]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} %{{.+}} + + // CK19: call void [[CALL22:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) + #pragma omp target map(tofrom:va[ii]) + { + va[15]++; + } + + // Always. + // Region 23 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE23]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE23]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8* + + // CK19: call void [[CALL23:@.+]](i32* {{[^,]+}}) + #pragma omp target map(always, tofrom: a) + { + a++; + } + + // Multidimensional arrays. + int marr[4][5][6]; + int ***mptr; + + // Region 24 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE24]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE24]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0]] to i8* + + // CK19: call void [[CALL24:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) + #pragma omp target map(tofrom: marr) + { + marr[1][2][3]++; + } + + // Region 25 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE25]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE25]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[6 x i32]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[SEC00]] = getelementptr {{.*}}[5 x [6 x i32]]* [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[SEC000]] = getelementptr {{.*}}[4 x [5 x [6 x i32]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + + // CK19: call void [[CALL25:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) + #pragma omp target map(tofrom: marr[1][2][2:4]) + { + marr[1][2][3]++; + } + + // Region 26 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE26]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE26]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[6 x i32]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[SEC00]] = getelementptr {{.*}}[5 x [6 x i32]]* [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[SEC000]] = getelementptr {{.*}}[4 x [5 x [6 x i32]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + + // CK19: call void [[CALL26:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) + #pragma omp target map(tofrom: marr[1][2][:]) + { + marr[1][2][3]++; + } + + // Region 27 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE27]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE27]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast [4 x [5 x [6 x i32]]]* [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[6 x i32]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 3 + // CK19-DAG: [[SEC00]] = getelementptr {{.*}}[5 x [6 x i32]]* [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[SEC000]] = getelementptr {{.*}}[4 x [5 x [6 x i32]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + + // CK19: call void [[CALL27:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) + #pragma omp target map(tofrom: marr[1][2][3]) + { + marr[1][2][3]++; + } + + // Region 28 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE28]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE28]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32*** [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32*** [[SEC0:%.+]] to i8* + // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]], + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]], + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32*** [[SEC0]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32** [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32** [[SEC11:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC11]] = load i32**, i32*** [[SEC111:%[^,]+]], + // CK19-DAG: [[SEC111]] = getelementptr {{.*}}i32*** [[SEC1111:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC1111]] = load i32***, i32**** [[PTR]], + + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: [[CBPVAL2]] = bitcast i32** [[SEC1]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast i32* [[SEC2:%.+]] to i8* + // CK19-DAG: [[SEC2]] = getelementptr {{.*}}i32* [[SEC22:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC22]] = load i32*, i32** [[SEC222:%[^,]+]], + // CK19-DAG: [[SEC222]] = getelementptr {{.*}}i32** [[SEC2222:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC2222]] = load i32**, i32*** [[SEC22222:%[^,]+]], + // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}i32*** [[SEC222222:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC222222]] = load i32***, i32**** [[PTR]], + + // CK19: call void [[CALL28:@.+]](i32*** {{[^,]+}}) + #pragma omp target map(tofrom: mptr[1][2][2:4]) + { + mptr[1][2][3]++; + } + + // Region 29 + // CK19-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[SIZE29]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE29]]{{.+}}) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK19-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK19-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK19-DAG: [[CBPVAL0]] = bitcast i32*** [[VAR0:%.+]] to i8* + // CK19-DAG: [[CPVAL0]] = bitcast i32*** [[SEC0:%.+]] to i8* + // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]], + // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]], + + // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK19-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK19-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK19-DAG: [[CBPVAL1]] = bitcast i32*** [[SEC0]] to i8* + // CK19-DAG: [[CPVAL1]] = bitcast i32** [[SEC1:%.+]] to i8* + // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32** [[SEC11:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC11]] = load i32**, i32*** [[SEC111:%[^,]+]], + // CK19-DAG: [[SEC111]] = getelementptr {{.*}}i32*** [[SEC1111:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC1111]] = load i32***, i32**** [[PTR]], + + // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK19-DAG: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] + // CK19-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] + // CK19-DAG: [[CBPVAL2]] = bitcast i32** [[SEC1]] to i8* + // CK19-DAG: [[CPVAL2]] = bitcast i32* [[SEC2:%.+]] to i8* + // CK19-DAG: [[SEC2]] = getelementptr {{.*}}i32* [[SEC22:[^,]+]], i{{.+}} 3 + // CK19-DAG: [[SEC22]] = load i32*, i32** [[SEC222:%[^,]+]], + // CK19-DAG: [[SEC222]] = getelementptr {{.*}}i32** [[SEC2222:[^,]+]], i{{.+}} 2 + // CK19-DAG: [[SEC2222]] = load i32**, i32*** [[SEC22222:%[^,]+]], + // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}i32*** [[SEC222222:[^,]+]], i{{.+}} 1 + // CK19-DAG: [[SEC222222]] = load i32***, i32**** [[PTR]], + + // CK19: call void [[CALL29:@.+]](i32*** {{[^,]+}}) + #pragma omp target map(tofrom: mptr[1][2][3]) + { + mptr[1][2][3]++; + } +} + +// CK19: define {{.+}}[[CALL00]] +// CK19: define {{.+}}[[CALL01]] +// CK19: define {{.+}}[[CALL02]] +// CK19: define {{.+}}[[CALL03]] +// CK19: define {{.+}}[[CALL04]] +// CK19: define {{.+}}[[CALL05]] +// CK19: define {{.+}}[[CALL06]] +// CK19: define {{.+}}[[CALL07]] +// CK19: define {{.+}}[[CALL08]] +// CK19: define {{.+}}[[CALL09]] +// CK19: define {{.+}}[[CALL10]] +// CK19: define {{.+}}[[CALL11]] +// CK19: define {{.+}}[[CALL12]] +// CK19: define {{.+}}[[CALL13]] +// CK19: define {{.+}}[[CALL14]] +// CK19: define {{.+}}[[CALL15]] +// CK19: define {{.+}}[[CALL16]] +// CK19: define {{.+}}[[CALL17]] +// CK19: define {{.+}}[[CALL18]] +// CK19: define {{.+}}[[CALL19]] +// CK19: define {{.+}}[[CALL20]] +// CK19: define {{.+}}[[CALL21]] +// CK19: define {{.+}}[[CALL22]] +// CK19: define {{.+}}[[CALL23]] +// CK19: define {{.+}}[[CALL24]] +// CK19: define {{.+}}[[CALL25]] +// CK19: define {{.+}}[[CALL26]] +// CK19: define {{.+}}[[CALL27]] +// CK19: define {{.+}}[[CALL28]] +// CK19: define {{.+}}[[CALL29]] + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK20 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-64 +// RUN: %clang_cc1 -DCK20 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK20 --check-prefix CK20-64 +// RUN: %clang_cc1 -DCK20 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-32 +// RUN: %clang_cc1 -DCK20 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK20 --check-prefix CK20-32 +#ifdef CK20 + +// CK20: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK20: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK20: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK20: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 1] + +// CK20: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK20: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK20: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 12] +// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK20-LABEL: explicit_maps_references_and_function_args +void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], float *d){ + + int &aa = a; + float &bb = b; + int (&cc)[10] = c; + float *&dd = d; + + // Region 00 + // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK20-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK20-DAG: [[CPVAL0]] = bitcast i32* [[RVAR00:%.+]] to i8* + // CK20-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK20-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK20: call void [[CALL00:@.+]](i32* {{[^,]+}}) + #pragma omp target map(to:aa) + { + aa += 1; + } + + // Region 01 + // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK20-DAG: [[CBPVAL0]] = bitcast [10 x i32]* [[RVAR0:%.+]] to i8* + // CK20-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK20-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[RVAR00:%.+]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[RVAR0]] = load [10 x i32]*, [10 x i32]** [[VAR0:%[^,]+]] + // CK20-DAG: [[RVAR00]] = load [10 x i32]*, [10 x i32]** [[VAR0]] + + // CK20: call void [[CALL01:@.+]]([10 x i32]* {{[^,]+}}) + #pragma omp target map(to:cc[:5]) + { + cc[3] += 1; + } + + // Region 02 + // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK20-DAG: [[CBPVAL0]] = bitcast float* [[VAR0:%.+]] to i8* + // CK20-DAG: [[CPVAL0]] = bitcast float* [[VAR0]] to i8* + + // CK20: call void [[CALL02:@.+]](float* {{[^,]+}}) + #pragma omp target map(from:b) + { + b += 1.0f; + } + + // Region 03 + // CK20-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK20-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK20-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK20-DAG: [[CBPVAL0]] = bitcast float* [[RVAR0:%.+]] to i8* + // CK20-DAG: [[CPVAL0]] = bitcast float* [[SEC0:%.+]] to i8* + // CK20-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]] + // CK20-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2 + // CK20-DAG: [[RVAR00]] = load float*, float** [[VAR0]] + + // CK20: call void [[CALL03:@.+]](float* {{[^,]+}}) + #pragma omp target map(from:d[2:3]) + { + d[2] += 1.0f; + } +} + +// CK20: define {{.+}}[[CALL00]] +// CK20: define {{.+}}[[CALL01]] +// CK20: define {{.+}}[[CALL02]] +// CK20: define {{.+}}[[CALL03]] + +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK21 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-64 +// RUN: %clang_cc1 -DCK21 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK21 --check-prefix CK21-64 +// RUN: %clang_cc1 -DCK21 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-32 +// RUN: %clang_cc1 -DCK21 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK21 --check-prefix CK21-32 +#ifdef CK21 +// CK21: [[ST:%.+]] = type { i32, i32, float* } + +// CK21: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK21: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK21: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492] +// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK21: [[SIZE02:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 500] +// CK21: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i32] [i32 2, i32 98] + +// CK21: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 492] +// CK21: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK21: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK21: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 2] + +// CK21: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] 4, i[[Z]] 4] +// CK21: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 67] + +// CK21-LABEL: explicit_maps_template_args_and_members + +template +struct CC { + T A; + int A2; + float *B; + + int foo(T arg) { + float la[X]; + T *lb; + + // Region 00 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0:%.+]], i{{.+}} 0, i{{.+}} 0 + + // CK21: call void [[CALL00:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(A) + { + A += 1; + } + + // Region 01 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK21-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] + // CK21-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 + // CK21-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] + + // CK21: call void [[CALL01:@.+]](i32* {{[^,]+}}) + #pragma omp target map(lb[:X]) + { + lb[4] += 1; + } + + // Region 02 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast float** [[SEC0:%.+]] to i8* + // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + + // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK21-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK21-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK21-DAG: [[CBPVAL1]] = bitcast float** [[SEC0]] to i8* + // CK21-DAG: [[CPVAL1]] = bitcast float* [[SEC1:%.+]] to i8* + // CK21-DAG: [[SEC1]] = getelementptr {{.*}}float* [[RVAR1:%[^,]+]], i{{.+}} 123 + // CK21-DAG: [[RVAR1]] = load float*, float** [[SEC1_:%[^,]+]] + // CK21-DAG: [[SEC1_]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 + + // CK21: call void [[CALL02:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(from:B[X:X+2]) + { + B[2] += 1.0f; + } + + // Region 03 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast [123 x float]* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast [123 x float]* [[VAR0]] to i8* + + // CK21: call void [[CALL03:@.+]]([123 x float]* {{[^,]+}}) + #pragma omp target map(from:la) + { + la[3] += 1.0f; + } + + // Region 04 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8* + + // CK21: call void [[CALL04:@.+]](i32* {{[^,]+}}) + #pragma omp target map(from:arg) + { + arg +=1; + } + + // Make sure the extra flag is passed to the second map. + // Region 05 + // CK21-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE05]]{{.+}}) + // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK21-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK21-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK21-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[VAR0:%.+]] to i8* + // CK21-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + + // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK21-DAG: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] + // CK21-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] + // CK21-DAG: [[CBPVAL1]] = bitcast [[ST]]* [[VAR1:%.+]] to i8* + // CK21-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* + // CK21-DAG: [[SEC1]] = getelementptr {{.*}}[[ST]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 + + // CK21: call void [[CALL05:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(A, A2) + { + A += 1; + A2 += 1; + } + return A; + } +}; + +int explicit_maps_template_args_and_members(int a){ + CC<123,int> c; + return c.foo(a); +} + +// CK21: define {{.+}}[[CALL00]] +// CK21: define {{.+}}[[CALL01]] +// CK21: define {{.+}}[[CALL02]] +// CK21: define {{.+}}[[CALL03]] +// CK21: define {{.+}}[[CALL04]] +// CK21: define {{.+}}[[CALL05]] +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK22 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-64 +// RUN: %clang_cc1 -DCK22 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK22 --check-prefix CK22-64 +// RUN: %clang_cc1 -DCK22 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-32 +// RUN: %clang_cc1 -DCK22 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK22 --check-prefix CK22-32 +#ifdef CK22 + +// CK22-DAG: [[ST:%.+]] = type { float } +// CK22-DAG: [[STT:%.+]] = type { i32 } + +// CK22: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK22: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK22: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK22: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK22: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK22: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK22: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE07:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK22: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE08:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK22: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK22: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE11:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK22: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE12:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK22: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK22: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK22: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +int a; +int c[100]; +int *d; + +struct ST { + float fa; +}; + +ST sa ; +ST sc[100]; +ST *sd; + +template +struct STT { + T fa; +}; + +STT sta ; +STT stc[100]; +STT *std; + +// CK22-LABEL: explicit_maps_globals +int explicit_maps_globals(void){ + // Region 00 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast (i32* @a to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast (i32* @a to i8*), i8** [[P0]] + + // CK22: call void [[CALL00:@.+]](i32* {{[^,]+}}) + #pragma omp target map(a) + { a+=1; } + + // Region 01 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x i32]* @c to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([100 x i32]* @c to i8*), i8** [[P0]] + + // CK22: call void [[CALL01:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(c) + { c[3]+=1; } + + // Region 02 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast (i32** @d to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast (i32** @d to i8*), i8** [[P0]] + + // CK22: call void [[CALL02:@.+]](i32** {{[^,]+}}) + #pragma omp target map(d) + { d[3]+=1; } + + // Region 03 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x i32]* @c to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast (i32* getelementptr inbounds ([100 x i32], [100 x i32]* @c, i{{.+}} 0, i{{.+}} 1) to i8*), i8** [[P0]] + + // CK22: call void [[CALL03:@.+]]([100 x i32]* {{[^,]+}}) + #pragma omp target map(c[1:4]) + { c[3]+=1; } + + // Region 04 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK22-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK22-DAG: [[CBPVAL0]] = bitcast i32* [[RVAR0:%.+]] to i8* + // CK22-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* + // CK22-DAG: [[RVAR0]] = load i32*, i32** @d + // CK22-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 2 + // CK22-DAG: [[RVAR00]] = load i32*, i32** @d + + // CK22: call void [[CALL04:@.+]](i32* {{[^,]+}}) + #pragma omp target map(d[2:5]) + { d[3]+=1; } + + // Region 05 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([[ST]]* @sa to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[ST]]* @sa to i8*), i8** [[P0]] + + // CK22: call void [[CALL05:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(sa) + { sa.fa+=1; } + + // Region 06 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE06]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x [[ST]]]* @sc to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([100 x [[ST]]]* @sc to i8*), i8** [[P0]] + + // CK22: call void [[CALL06:@.+]]([100 x [[ST]]]* {{[^,]+}}) + #pragma omp target map(sc) + { sc[3].fa+=1; } + + // Region 07 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE07]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE07]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([[ST]]** @sd to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[ST]]** @sd to i8*), i8** [[P0]] + + // CK22: call void [[CALL07:@.+]]([[ST]]** {{[^,]+}}) + #pragma omp target map(sd) + { sd[3].fa+=1; } + + // Region 08 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE08]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x [[ST]]]* @sc to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[ST]]* getelementptr inbounds ([100 x [[ST]]], [100 x [[ST]]]* @sc, i{{.+}} 0, i{{.+}} 1) to i8*), i8** [[P0]] + + // CK22: call void [[CALL08:@.+]]([100 x [[ST]]]* {{[^,]+}}) + #pragma omp target map(sc[1:4]) + { sc[3].fa+=1; } + + // Region 09 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK22-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK22-DAG: [[CBPVAL0]] = bitcast [[ST]]* [[RVAR0:%.+]] to i8* + // CK22-DAG: [[CPVAL0]] = bitcast [[ST]]* [[SEC0:%.+]] to i8* + // CK22-DAG: [[RVAR0]] = load [[ST]]*, [[ST]]** @sd + // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[RVAR00:%.+]], i{{.+}} 2 + // CK22-DAG: [[RVAR00]] = load [[ST]]*, [[ST]]** @sd + + // CK22: call void [[CALL09:@.+]]([[ST]]* {{[^,]+}}) + #pragma omp target map(sd[2:5]) + { sd[3].fa+=1; } + + // Region 10 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([[STT]]* @sta to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[STT]]* @sta to i8*), i8** [[P0]] + + // CK22: call void [[CALL10:@.+]]([[STT]]* {{[^,]+}}) + #pragma omp target map(sta) + { sta.fa+=1; } + + // Region 11 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE11]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x [[STT]]]* @stc to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([100 x [[STT]]]* @stc to i8*), i8** [[P0]] + + // CK22: call void [[CALL11:@.+]]([100 x [[STT]]]* {{[^,]+}}) + #pragma omp target map(stc) + { stc[3].fa+=1; } + + // Region 12 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE12]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([[STT]]** @std to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[STT]]** @std to i8*), i8** [[P0]] + + // CK22: call void [[CALL12:@.+]]([[STT]]** {{[^,]+}}) + #pragma omp target map(std) + { std[3].fa+=1; } + + // Region 13 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE13]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* bitcast ([100 x [[STT]]]* @stc to i8*), i8** [[BP0]] + // CK22-DAG: store i8* bitcast ([[STT]]* getelementptr inbounds ([100 x [[STT]]], [100 x [[STT]]]* @stc, i{{.+}} 0, i{{.+}} 1) to i8*), i8** [[P0]] + + // CK22: call void [[CALL13:@.+]]([100 x [[STT]]]* {{[^,]+}}) + #pragma omp target map(stc[1:4]) + { stc[3].fa+=1; } + + // Region 14 + // CK22-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE14]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}}) + // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK22-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK22-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK22-DAG: [[CBPVAL0]] = bitcast [[STT]]* [[RVAR0:%.+]] to i8* + // CK22-DAG: [[CPVAL0]] = bitcast [[STT]]* [[SEC0:%.+]] to i8* + // CK22-DAG: [[RVAR0]] = load [[STT]]*, [[STT]]** @std + // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[STT]]* [[RVAR00:%.+]], i{{.+}} 2 + // CK22-DAG: [[RVAR00]] = load [[STT]]*, [[STT]]** @std + + // CK22: call void [[CALL14:@.+]]([[STT]]* {{[^,]+}}) + #pragma omp target map(std[2:5]) + { std[3].fa+=1; } + + return 0; +} +// CK22: define {{.+}}[[CALL00]] +// CK22: define {{.+}}[[CALL01]] +// CK22: define {{.+}}[[CALL02]] +// CK22: define {{.+}}[[CALL03]] +// CK22: define {{.+}}[[CALL04]] +// CK22: define {{.+}}[[CALL05]] +// CK22: define {{.+}}[[CALL06]] +// CK22: define {{.+}}[[CALL07]] +// CK22: define {{.+}}[[CALL08]] +// CK22: define {{.+}}[[CALL09]] +// CK22: define {{.+}}[[CALL10]] +// CK22: define {{.+}}[[CALL11]] +// CK22: define {{.+}}[[CALL12]] +// CK22: define {{.+}}[[CALL13]] +// CK22: define {{.+}}[[CALL14]] +#endif +///==========================================================================/// +// RUN: %clang_cc1 -std=c++11 -DCK23 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64 +// RUN: %clang_cc1 -std=c++11 -DCK23 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -std=c++11 -fopenmp -omptargets=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 CK23 --check-prefix CK23-64 +// RUN: %clang_cc1 -std=c++11 -DCK23 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-32 +// RUN: %clang_cc1 -std=c++11 -DCK23 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -std=c++11 -fopenmp -omptargets=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 CK23 --check-prefix CK23-32 +#ifdef CK23 + +// CK23: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK23: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK23: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400] +// CK23: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK23: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK23: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 16] +// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK23-LABEL: explicit_maps_inside_captured +int explicit_maps_inside_captured(int a){ + float b; + float c[100]; + float *d; + + // CK23: call void @{{.*}}explicit_maps_inside_captured{{.*}}([[SA:%.+]]* {{.*}}) + // CK23: define {{.*}}explicit_maps_inside_captured{{.*}} + [&](void){ + // Region 00 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast i32* [[VAR00:%.+]] to i8* + // CK23-DAG: [[VAR0]] = load i32*, i32** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load i32*, i32** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL00:@.+]](i32* {{[^,]+}}) + #pragma omp target map(a) + { a+=1; } + // Region 01 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast float* [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast float* [[VAR00:%.+]] to i8* + // CK23-DAG: [[VAR0]] = load float*, float** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load float*, float** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL01:@.+]](float* {{[^,]+}}) + #pragma omp target map(b) + { b+=1; } + // Region 02 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast [100 x float]* [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast [100 x float]* [[VAR00:%.+]] to i8* + // CK23-DAG: [[VAR0]] = load [100 x float]*, [100 x float]** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load [100 x float]*, [100 x float]** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL02:@.+]]([100 x float]* {{[^,]+}}) + #pragma omp target map(c) + { c[3]+=1; } + + // Region 03 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast float** [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast float** [[VAR00:%.+]] to i8* + // CK23-DAG: [[VAR0]] = load float**, float*** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load float**, float*** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL03:@.+]](float** {{[^,]+}}) + #pragma omp target map(d) + { d[3]+=1; } + // Region 04 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast [100 x float]* [[VAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast float* [[SEC0:%.+]] to i8* + // CK23-DAG: [[SEC0]] = getelementptr {{.*}}[100 x float]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2 + // CK23-DAG: [[VAR0]] = load [100 x float]*, [100 x float]** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load [100 x float]*, [100 x float]** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL04:@.+]]([100 x float]* {{[^,]+}}) + #pragma omp target map(c[2:4]) + { c[3]+=1; } + + // Region 05 + // CK23-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}) + // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + + // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK23-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] + // CK23-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] + // CK23-DAG: [[CBPVAL0]] = bitcast float* [[RVAR0:%.+]] to i8* + // CK23-DAG: [[CPVAL0]] = bitcast float* [[SEC0:%.+]] to i8* + // CK23-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]] + // CK23-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2 + // CK23-DAG: [[RVAR00]] = load float*, float** [[VAR00:%[^,]+]] + // CK23-DAG: [[VAR0]] = load float**, float*** [[CAP0:%[^,]+]] + // CK23-DAG: [[CAP0]] = getelementptr inbounds [[SA]], [[SA]] + // CK23-DAG: [[VAR00]] = load float**, float*** [[CAP00:%[^,]+]] + // CK23-DAG: [[CAP00]] = getelementptr inbounds [[SA]], [[SA]] + + // CK23: call void [[CALL05:@.+]](float* {{[^,]+}}) + #pragma omp target map(d[2:4]) + { d[3]+=1; } + }(); + return b; +} + +// CK23: define {{.+}}[[CALL00]] +// CK23: define {{.+}}[[CALL01]] +// CK23: define {{.+}}[[CALL02]] +// CK23: define {{.+}}[[CALL03]] +// CK23: define {{.+}}[[CALL04]] +// CK23: define {{.+}}[[CALL05]] +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK24 -verify -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-64 +// RUN: %clang_cc1 -DCK24 -fopenmp -omptargets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK24 --check-prefix CK24-64 +// RUN: %clang_cc1 -DCK24 -verify -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-32 +// RUN: %clang_cc1 -DCK24 -fopenmp -omptargets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -omptargets=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 CK24 --check-prefix CK24-32 +#ifdef CK24 + +// CK24-DAG: [[SC:%.+]] = type { i32, [[SB:%.+]], [[SB:%.+]]*, [10 x i32] } +// CK24-DAG: [[SB]] = type { i32, [[SA:%.+]], [10 x [[SA:%.+]]], [10 x [[SA:%.+]]*], [[SA:%.+]]* } +// CK24-DAG: [[SA]] = type { i32, [[SA]]*, [10 x i32] } + +struct SA{ + int a; + struct SA *p; + int b[10]; +}; +struct SB{ + int a; + struct SA s; + struct SA sa[10]; + struct SA *sp[10]; + struct SA *p; +}; +struct SC{ + int a; + struct SB s; + struct SB *p; + int b[10]; +}; + +// CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] +// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}] +// CK24: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK24: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}] +// CK24: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE07:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE07:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE08:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE08:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE09:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE09:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 8] +// CK24: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE11:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}] +// CK24: [[MTYPE11:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// 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 3, i32 99, i32 99, i32 99] + +// CK24: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE14:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}] +// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE15:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE16:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] +// CK24: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE17:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}] +// CK24: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE18:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] +// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE19:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE20:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE21:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE22:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{8|4}}] +// CK24: [[MTYPE22:@.+]] = private {{.*}}constant [1 x i32] [i32 3] + +// CK24: [[SIZE23:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}] +// CK24: [[MTYPE23:@.+]] = private {{.*}}constant [2 x i32] [i32 3, i32 99] + +// CK24: [[SIZE24:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4] +// CK24: [[MTYPE24:@.+]] = private {{.*}}constant [4 x i32] [i32 3, i32 99, i32 99, i32 99] + +// CK24-LABEL: explicit_maps_struct_fields +int explicit_maps_struct_fields(int a){ + SC s; + SC *p; + +// Region 01 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 + +// CK24: call void [[CALL01:@.+]]([[SC]]* {{[^,]+}}) +#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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]* [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast [[SB]]* [[SEC1:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast [[SA]]** [[SEC1:%.+]] to i8* +// 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: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] +// CK24-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] +// CK24-DAG: [[CBPVAL2]] = bitcast [[SA]]** [[SEC1]] to i8* +// CK24-DAG: [[CPVAL2]] = bitcast [[SA]]** [[SEC2:%.+]] to i8* +// 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: store i8* [[CBPVAL3:%[^,]+]], i8** [[BP3]] +// CK24-DAG: store i8* [[CPVAL3:%[^,]+]], i8** [[P3]] +// CK24-DAG: [[CBPVAL3]] = bitcast [[SA]]** [[SEC2]] to i8* +// CK24-DAG: [[CPVAL3]] = bitcast i32* [[SEC3:%.+]] to i8* +// 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. +// +// Region 13 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE13]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 0 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL13:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->a) + { p->a++; } + +// Region 14 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE14]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL14:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.s) + { p->a++; } + +// Region 15 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE15]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE15]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// 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]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL15:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.s.a) + { p->a++; } + +// Region 16 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE16]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE16]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 3 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL16:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->b[:5]) + { p->a++; } + +// Region 17 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE17]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE17]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast [[SB]]* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL17:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->p[:5]) + { p->a++; } + +// Region 18 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE18]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE18]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// 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]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL18:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.sa[3].a) + { p->a++; } + +// Region 19 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE19]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE19]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// 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]]* [[VAR00:%.+]], 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// 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]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL19:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.sp[3]->a) + { p->a++; } + +// Region 20 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE20]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE20]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0 +// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL20:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->p->a) + { p->a++; } + +// Region 21 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE21]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE21]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// 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]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL21:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.p->a) + { p->a++; } + +// Region 22 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE22]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE22]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast i32* [[SEC0:%.+]] to i8* +// 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]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL22:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.s.b[:2]) + { p->a++; } + +// Region 23 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE23]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE23]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SA]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SA]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast i32* [[SEC1:%.+]] to i8* +// 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]]* [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL23:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->s.p->b[:2]) + { p->a++; } + +// Region 24 +// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[SIZE24]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE24]]{{.+}}) +// 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: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]] +// CK24-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]] +// CK24-DAG: [[CBPVAL0]] = bitcast [[SC]]* [[VAR0:%.+]] to i8* +// CK24-DAG: [[CPVAL0]] = bitcast [[SB]]** [[SEC0:%.+]] to i8* +// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], 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: store i8* [[CBPVAL1:%[^,]+]], i8** [[BP1]] +// CK24-DAG: store i8* [[CPVAL1:%[^,]+]], i8** [[P1]] +// CK24-DAG: [[CBPVAL1]] = bitcast [[SB]]** [[SEC0]] to i8* +// CK24-DAG: [[CPVAL1]] = bitcast [[SA]]** [[SEC1:%.+]] to i8* +// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 4 +// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], +// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR000:%.+]], 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: store i8* [[CBPVAL2:%[^,]+]], i8** [[BP2]] +// CK24-DAG: store i8* [[CPVAL2:%[^,]+]], i8** [[P2]] +// CK24-DAG: [[CBPVAL2]] = bitcast [[SA]]** [[SEC1]] to i8* +// CK24-DAG: [[CPVAL2]] = bitcast [[SA]]** [[SEC2:%.+]] to i8* +// 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]]* [[VAR0000:%.+]], 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: store i8* [[CBPVAL3:%[^,]+]], i8** [[BP3]] +// CK24-DAG: store i8* [[CPVAL3:%[^,]+]], i8** [[P3]] +// CK24-DAG: [[CBPVAL3]] = bitcast [[SA]]** [[SEC2]] to i8* +// CK24-DAG: [[CPVAL3]] = bitcast i32* [[SEC3:%.+]] to i8* +// 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]]* [[VAR00000:%.+]], i{{.+}} 0, i{{.+}} 2 + +// CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR000]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR0000]] = load [[SC]]*, [[SC]]** %{{.+}} +// CK24-DAG: [[VAR00000]] = load [[SC]]*, [[SC]]** %{{.+}} + +// CK24: call void [[CALL24:@.+]]([[SC]]* {{[^,]+}}) +#pragma omp target map(p->p->p->p->a) + { p->a++; } + + return s.a; +} + +// 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]] +// CK24: define {{.+}}[[CALL16]] +// CK24: define {{.+}}[[CALL17]] +// CK24: define {{.+}}[[CALL18]] +// CK24: define {{.+}}[[CALL19]] +// CK24: define {{.+}}[[CALL20]] +// CK24: define {{.+}}[[CALL21]] +// CK24: define {{.+}}[[CALL22]] +// CK24: define {{.+}}[[CALL23]] +// CK24: define {{.+}}[[CALL24]] +#endif + #endif