Index: include/clang/AST/GlobalDecl.h =================================================================== --- include/clang/AST/GlobalDecl.h +++ include/clang/AST/GlobalDecl.h @@ -53,6 +53,7 @@ GlobalDecl(const CapturedDecl *D) { Init(D); } GlobalDecl(const ObjCMethodDecl *D) { Init(D); } GlobalDecl(const OMPDeclareReductionDecl *D) { Init(D); } + GlobalDecl(const OMPDeclareMapperDecl *D) { Init(D); } GlobalDecl(const CXXConstructorDecl *D, CXXCtorType Type) : Value(D, Type) {} GlobalDecl(const CXXDestructorDecl *D, CXXDtorType Type) : Value(D, Type) {} Index: lib/AST/ASTContext.cpp =================================================================== --- lib/AST/ASTContext.cpp +++ lib/AST/ASTContext.cpp @@ -9804,7 +9804,7 @@ return !D->getDeclContext()->isDependentContext(); else if (isa(D)) return !D->getDeclContext()->isDependentContext(); - else if (isa(D)) + else if (isa(D) || isa(D)) return !D->getDeclContext()->isDependentContext(); else if (isa(D)) return true; Index: lib/CodeGen/CGDecl.cpp =================================================================== --- lib/CodeGen/CGDecl.cpp +++ lib/CodeGen/CGDecl.cpp @@ -143,7 +143,7 @@ return CGM.EmitOMPDeclareReduction(cast(&D), this); case Decl::OMPDeclareMapper: - return CGM.EmitOMPDeclareMapper(cast(&D), this); + return CGM.EmitOMPDeclareMapper(cast(&D)); case Decl::Typedef: // typedef int X; case Decl::TypeAlias: { // using X = int; [C++0x] @@ -2566,11 +2566,11 @@ getOpenMPRuntime().emitUserDefinedReduction(CGF, D); } -void CodeGenModule::EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D, - CodeGenFunction *CGF) { - if (!LangOpts.OpenMP || (!LangOpts.EmitAllDecls && !D->isUsed())) +void CodeGenModule::EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D) { + if (!LangOpts.OpenMP || LangOpts.OpenMPSimd || + (!LangOpts.EmitAllDecls && !D->isUsed())) return; - // FIXME: need to implement mapper code generation + getOpenMPRuntime().emitUserDefinedMapper(D); } void CodeGenModule::EmitOMPRequiresDecl(const OMPRequiresDecl *D) { Index: lib/CodeGen/CGOpenMPRuntime.h =================================================================== --- lib/CodeGen/CGOpenMPRuntime.h +++ lib/CodeGen/CGOpenMPRuntime.h @@ -345,6 +345,12 @@ SmallVector> FunctionUDRMapTy; FunctionUDRMapTy FunctionUDRMap; + /// Map from the user-defined mapper declaration to its corresponding + /// functions. The first one is the synchronous version, while the second one + /// is the asynchronous version. + llvm::DenseMap> + UDMMap; /// Type kmp_critical_name, originally defined as typedef kmp_int32 /// kmp_critical_name[8]; llvm::ArrayType *KmpCriticalNameTy; @@ -782,6 +788,22 @@ virtual std::pair getUserDefinedReduction(const OMPDeclareReductionDecl *D); + /// Emit code for the user defined mapper construct. + virtual void emitUserDefinedMapper(const OMPDeclareMapperDecl *D); + + /// Emit a function for a user defined mapper. Whether it is synchronous or + /// asynchronous depends on \a NoWait. + virtual llvm::Function *emitUDMapperFunc(const OMPDeclareMapperDecl *D, + bool NoWait); + + // Emit the array initialization or deletion portion for user-defined mapper + // code generation. + virtual llvm::Value * + emitUDMapperArrayInitOrDel(CodeGenFunction &MapperCGF, llvm::Value *DeviceID, + llvm::Value *BasePtr, llvm::Value *Ptr, + llvm::Value *Size, llvm::Value *MapType, + CharUnits ElementSize, bool IsInit, bool NoWait); + /// Emits outlined function for the specified OpenMP parallel directive /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID, /// kmp_int32 BoundID, struct context_vars*). @@ -2060,6 +2082,9 @@ bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) override; + /// Emit code for the user defined mapper construct. + void emitUserDefinedMapper(const OMPDeclareMapperDecl *D); + /// Emit the target offloading code associated with \a D. The emitted /// code attempts offloading the execution to the device, an the event of /// a failure it executes the host version outlined in \a OutlinedFn. Index: lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenMPRuntime.cpp +++ lib/CodeGen/CGOpenMPRuntime.cpp @@ -715,6 +715,14 @@ // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t // *arg_types); OMPRTL__tgt_target_data_update_nowait, + // Call to void __tgt_target_data_mapper(int64_t device_id, int32_t arg_num, + // void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types, void + // **arg_mapper_ptrs); + OMPRTL__tgt_target_data_mapper, + // Call to void __tgt_target_data_mapper_nowait(int64_t device_id, int32_t + // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t + // *arg_types, void **arg_mapper_ptrs); + OMPRTL__tgt_target_data_mapper_nowait, }; /// A basic class for pre|post-action for advanced codegen sequence for OpenMP @@ -2379,6 +2387,38 @@ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_nowait"); break; } + case OMPRTL__tgt_target_data_mapper: { + // Build void __tgt_target_data_mapper(int64_t device_id, int32_t arg_num, + // void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types, + // void **arg_mapper_ptrs); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.SizeTy->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.IntTy, TypeParams, /*isVarArg*/ false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_mapper"); + break; + } + case OMPRTL__tgt_target_data_mapper_nowait: { + // Build void __tgt_target_data_mapper_nowait(int64_t device_id, int32_t + // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t + // *arg_types, void **arg_mapper_ptrs); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.SizeTy->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.IntTy, TypeParams, /*isVarArg*/ false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_mapper_nowait"); + break; + } } assert(RTLFn && "Unable to find OpenMP runtime function"); return RTLFn; @@ -6724,8 +6764,13 @@ : IE(IE), VD(VD) {} }; - /// Directive from where the map clauses were extracted. - const OMPExecutableDirective &CurDir; + /// The target directive from where the mappable clauses were extracted. One + /// and only one of CurDir and CurMapperDir is valid. + const OMPExecutableDirective *CurDir = nullptr; + + /// The mapper directive from where the map clauses were extracted. One and + /// only one of CurDir and CurMapperDir is valid. + const OMPDeclareMapperDecl *CurMapperDir = nullptr; /// Function the directive is being generated for. CodeGenFunction &CGF; @@ -7395,7 +7440,7 @@ public: MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF) - : CurDir(Dir), CGF(CGF) { + : CurDir(&Dir), CGF(CGF) { // Extract firstprivate clause information. for (const auto *C : Dir.getClausesOfKind()) for (const auto *D : C->varlists()) @@ -7407,6 +7452,10 @@ DevPointersMap[L.first].push_back(L.second); } + /// Constructor for the declare mapper directive. + MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF) + : CurMapperDir(&Dir), CGF(CGF) {} + /// Generate code for the combined entry if we have a partially mapped struct /// and take care of the mapping flags of the arguments corresponding to /// individual struct members. @@ -7469,17 +7518,17 @@ }; // FIXME: MSVC 2013 seems to require this-> to find member CurDir. - for (const auto *C : this->CurDir.getClausesOfKind()) + for (const auto *C : this->CurDir->getClausesOfKind()) for (const auto &L : C->component_lists()) { InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifiers(), /*ReturnDevicePointer=*/false, C->isImplicit()); } - for (const auto *C : this->CurDir.getClausesOfKind()) + for (const auto *C : this->CurDir->getClausesOfKind()) for (const auto &L : C->component_lists()) { InfoGen(L.first, L.second, OMPC_MAP_to, llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit()); } - for (const auto *C : this->CurDir.getClausesOfKind()) + for (const auto *C : this->CurDir->getClausesOfKind()) for (const auto &L : C->component_lists()) { InfoGen(L.first, L.second, OMPC_MAP_from, llvm::None, /*ReturnDevicePointer=*/false, C->isImplicit()); @@ -7496,7 +7545,7 @@ // FIXME: MSVC 2013 seems to require this-> to find member CurDir. for (const auto *C : - this->CurDir.getClausesOfKind()) { + this->CurDir->getClausesOfKind()) { for (const auto &L : C->component_lists()) { assert(!L.second.empty() && "Not expecting empty list of components!"); const ValueDecl *VD = L.second.back().getAssociatedDeclaration(); @@ -7622,6 +7671,80 @@ } } + /// Generate all the base pointers, section pointers, sizes and map types for + /// the extracted map clauses. + void generateAllInfoForMapper(MapBaseValuesArrayTy &BasePointers, + MapValuesArrayTy &Pointers, + MapValuesArrayTy &Sizes, + MapFlagsArrayTy &Types) const { + // FIXME: MSVC 2013 seems to require this-> to find members. + assert(this->CurMapperDir && !this->CurDir && + "Expect a valid declare mapper directive."); + // We have to process the component lists that relate with the same + // declaration in a single chunk so that we can generate the map flags + // correctly. Therefore, we organize all lists in a map. + llvm::MapVector> Info; + + // Helper function to fill the information map for the different supported + // clauses. + auto &&InfoGen = [&Info]( + const ValueDecl *D, + OMPClauseMappableExprCommon::MappableExprComponentListRef L, + OpenMPMapClauseKind MapType, + ArrayRef MapModifiers, + bool ReturnDevicePointer, bool IsImplicit) { + const ValueDecl *VD = + D ? cast(D->getCanonicalDecl()) : nullptr; + Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer, + IsImplicit); + }; + + // FIXME: MSVC 2013 seems to require this-> to find member CurMapperDir. + for (const auto *C : this->CurMapperDir->clauselists()) { + const auto *MC = cast(C); + for (const auto &L : MC->component_lists()) { + InfoGen(L.first, L.second, MC->getMapType(), MC->getMapTypeModifiers(), + /*ReturnDevicePointer=*/false, MC->isImplicit()); + } + } + + for (const auto &M : Info) { + // We need to know when we generate information for the first component + // associated with a capture, because the mapping flags depend on it. + bool IsFirstComponentList = true; + + // Temporary versions of arrays + MapBaseValuesArrayTy CurBasePointers; + MapValuesArrayTy CurPointers; + MapValuesArrayTy CurSizes; + MapFlagsArrayTy CurTypes; + StructRangeInfoTy PartialStruct; + + for (const MapInfo &L : M.second) { + assert(!L.Components.empty() && + "Not expecting declaration with no component lists."); + // FIXME: MSVC 2013 seems to require this-> to find the member method. + this->generateInfoForComponentList( + L.MapType, L.MapModifiers, L.Components, CurBasePointers, + CurPointers, CurSizes, CurTypes, PartialStruct, + IsFirstComponentList, L.IsImplicit); + IsFirstComponentList = false; + } + + // If there is an entry in PartialStruct it means we have a struct with + // individual members mapped. Emit an extra combined entry. + if (PartialStruct.Base.isValid()) + emitCombinedEntry(BasePointers, Pointers, Sizes, Types, CurTypes, + PartialStruct); + + // 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()); + Types.append(CurTypes.begin(), CurTypes.end()); + } + } + /// Emit capture info for lambdas for variables captured by reference. void generateInfoForLambdaCaptures( const ValueDecl *VD, llvm::Value *Arg, MapBaseValuesArrayTy &BasePointers, @@ -7730,7 +7853,7 @@ OpenMPMapClauseKind, ArrayRef, bool>; SmallVector DeclComponentLists; // FIXME: MSVC 2013 seems to require this-> to find member CurDir. - for (const auto *C : this->CurDir.getClausesOfKind()) { + for (const auto *C : this->CurDir->getClausesOfKind()) { for (const auto &L : C->decl_component_lists(VD)) { assert(L.first == VD && "We got information for the wrong declaration??"); @@ -7880,7 +8003,7 @@ MapFlagsArrayTy &Types) const { // Map other list items in the map clause which are not captured variables // but "declare target link" global variables., - for (const auto *C : this->CurDir.getClausesOfKind()) { + for (const auto *C : this->CurDir->getClausesOfKind()) { for (const auto &L : C->component_lists()) { if (!L.first) continue; @@ -7984,7 +8107,8 @@ MappableExprsHandler::MapValuesArrayTy &Pointers, MappableExprsHandler::MapValuesArrayTy &Sizes, MappableExprsHandler::MapFlagsArrayTy &MapTypes, - CGOpenMPRuntime::TargetDataInfo &Info) { + CGOpenMPRuntime::TargetDataInfo &Info, + llvm::Value *MapperMapType = nullptr) { CodeGenModule &CGM = CGF.CGM; ASTContext &Ctx = CGF.getContext(); @@ -8002,6 +8126,9 @@ break; } + // Indicate whether it is code generation within a user-defined mapper. + bool IsMapper = MapperMapType; + llvm::APInt PointerNumAP(32, Info.NumberOfPtrs, /*isSigned=*/true); QualType PointerArrayType = Ctx.getConstantArrayType(Ctx.VoidPtrTy, PointerNumAP, ArrayType::Normal, @@ -8039,20 +8166,31 @@ Info.SizesArray = SizesArrayGbl; } - // The map types are always constant so we don't need to generate code to - // fill arrays. Instead, we create an array constant. - SmallVector Mapping(MapTypes.size(), 0); - llvm::copy(MapTypes, Mapping.begin()); - llvm::Constant *MapTypesArrayInit = - llvm::ConstantDataArray::get(CGF.Builder.getContext(), Mapping); - std::string MaptypesName = - CGM.getOpenMPRuntime().getName({"offload_maptypes"}); - auto *MapTypesArrayGbl = new llvm::GlobalVariable( - CGM.getModule(), MapTypesArrayInit->getType(), - /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage, - MapTypesArrayInit, MaptypesName); - MapTypesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global); - Info.MapTypesArray = MapTypesArrayGbl; + if (IsMapper) { + // Allocate the memory space for map types because they are decided + // dynamically. + QualType MapArrayType = Ctx.getConstantArrayType( + Ctx.getIntTypeForBitwidth(/*DestWidth*/ 64, /*Signed*/ true), + PointerNumAP, ArrayType::Normal, + /*IndexTypeQuals=*/0); + Info.MapTypesArray = + CGF.CreateMemTemp(MapArrayType, ".offload_maptypes").getPointer(); + } else { + // The map types are always constant so we don't need to generate code to + // fill arrays. Instead, we create an array constant. + SmallVector Mapping(MapTypes.size(), 0); + llvm::copy(MapTypes, Mapping.begin()); + llvm::Constant *MapTypesArrayInit = + llvm::ConstantDataArray::get(CGF.Builder.getContext(), Mapping); + std::string MaptypesName = + CGM.getOpenMPRuntime().getName({"offload_maptypes"}); + auto *MapTypesArrayGbl = new llvm::GlobalVariable( + CGM.getModule(), MapTypesArrayInit->getType(), + /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage, + MapTypesArrayInit, MaptypesName); + MapTypesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global); + Info.MapTypesArray = MapTypesArrayGbl; + } for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) { llvm::Value *BPVal = *BasePointers[I]; @@ -8088,9 +8226,77 @@ CGF.Builder.CreateIntCast(Sizes[I], CGM.SizeTy, /*isSigned=*/true), SAddr); } + + if (IsMapper) { + // Combine the map type inherited from user-defined mapper with that + // specified in the program. + // [OpenMP 5.0], 1.2.6. map-type decay. + // | alloc | to | from | tofrom | release | delete + // ---------------------------------------------------------- + // alloc | alloc | alloc | alloc | alloc | release | delete + // to | alloc | to | alloc | to | release | delete + // from | alloc | alloc | from | from | release | delete + // tofrom | alloc | to | from | tofrom | release | delete + llvm::Value *GEP = CGF.Builder.CreateConstInBoundsGEP2_32( + llvm::ArrayType::get(CGM.Int64Ty, Info.NumberOfPtrs), + Info.MapTypesArray, /*Idx0=*/0, /*Idx1=*/I); + llvm::Value *OriMapType = CGF.Builder.getInt64(MapTypes[I]); + llvm::Value *LeftToFrom = CGF.Builder.CreateAnd( + MapperMapType, + CGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_TO | + MappableExprsHandler::OMP_MAP_FROM)); + llvm::BasicBlock *AllocBB = CGF.createBasicBlock("omp.type.alloc"); + llvm::BasicBlock *AllocElseBB = + CGF.createBasicBlock("omp.type.alloc.else"); + llvm::BasicBlock *ToBB = CGF.createBasicBlock("omp.type.to"); + llvm::BasicBlock *ToElseBB = CGF.createBasicBlock("omp.type.to.else"); + llvm::BasicBlock *FromBB = CGF.createBasicBlock("omp.type.from"); + llvm::BasicBlock *EndBB = CGF.createBasicBlock("omp.type.end"); + llvm::Value *IsAlloc = CGF.Builder.CreateIsNull(LeftToFrom); + CGF.Builder.CreateCondBr(IsAlloc, AllocBB, AllocElseBB); + // In case of alloc, clear OMP_MAP_TO and OMP_MAP_FROM. + CGF.EmitBlock(AllocBB); + llvm::Value *AllocMapType = CGF.Builder.CreateAnd( + OriMapType, + CGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO | + MappableExprsHandler::OMP_MAP_FROM))); + CGF.Builder.CreateBr(EndBB); + CGF.EmitBlock(AllocElseBB); + llvm::Value *IsTo = CGF.Builder.CreateICmpEQ( + LeftToFrom, CGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_TO)); + CGF.Builder.CreateCondBr(IsTo, ToBB, ToElseBB); + // In case of to, clear OMP_MAP_FROM. + CGF.EmitBlock(ToBB); + llvm::Value *ToMapType = CGF.Builder.CreateAnd( + OriMapType, + CGF.Builder.getInt64(~MappableExprsHandler::OMP_MAP_FROM)); + CGF.Builder.CreateBr(EndBB); + CGF.EmitBlock(ToElseBB); + llvm::Value *IsFrom = CGF.Builder.CreateICmpEQ( + LeftToFrom, + CGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_FROM)); + CGF.Builder.CreateCondBr(IsFrom, FromBB, EndBB); + // In case of from, clear OMP_MAP_TO. + CGF.EmitBlock(FromBB); + llvm::Value *FromMapType = CGF.Builder.CreateAnd( + OriMapType, + CGF.Builder.getInt64(~MappableExprsHandler::OMP_MAP_TO)); + // In case of tofrom, do nothing. + CGF.EmitBlock(EndBB); + llvm::PHINode *MapType = + CGF.Builder.CreatePHI(CGM.Int64Ty, 4, "omp.maptype"); + MapType->addIncoming(AllocMapType, AllocBB); + MapType->addIncoming(ToMapType, ToBB); + MapType->addIncoming(FromMapType, FromBB); + MapType->addIncoming(OriMapType, ToElseBB); + Address Addr(GEP, Ctx.getTypeAlignInChars(Ctx.getIntTypeForBitwidth( + /*DestWidth*/ 64, /*Signed*/ true))); + CGF.Builder.CreateStore(MapType, Addr); + } } } } + /// Emit the arguments to be passed to the runtime library based on the /// arrays of pointers, sizes and map types. static void emitOffloadingArraysArgument( @@ -8273,6 +8479,322 @@ return nullptr; } +/// Emit code for the user defined mapper construct. +void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D) { + if (UDMMap.count(D) > 0) + return; + // Generate a synchronous mapper function. + llvm::Function *SyncFn = emitUDMapperFunc(D, /*NoWait=*/false); + // Generate an asynchronous mapper function. + llvm::Function *AsyncFn = emitUDMapperFunc(D, /*NoWait=*/true); + // Add the generated mapper functions to UDMMap. + UDMMap.try_emplace(D, SyncFn, AsyncFn); +} + +/// Emit the user-defined mapper function. Whether it is synchronous or +/// asynchronous depends on \a NoWait. The code generation follows the pattern +/// in the example below. +/// \code +/// int .omp_mapper_.(int64_t device_id, Ty *base_ptr, Ty *ptr, +/// size_t size, int64_t maptype) { +/// // Allocate space for an array section first. +/// if (size > 1 && !maptype.IsDelete) { +/// res = __tgt_target_data_mapper(device_id, /*arg_num*/1, &base_ptr, &ptr, +/// size*sizeof(Ty), maptype, null); +/// if (res != 0) // Data mapping failed. +/// return res; +/// } +/// // Map members. +/// for (unsigned i = 0; i < size; i++) { +/// ...; // Prepare arguments of __tgt_target_data_mapper. +/// res = __tgt_target_data_mapper(device_id, arg_num, arg_base, arg, size, +/// maptype, mapper); +/// if (res != 0) // Data mapping failed. +/// return res; +/// } +/// // Delete the array section. +/// if (size > 1 && maptype.IsDelete) { +/// res = __tgt_target_data_mapper(device_id, /*arg_num*/1, &base_ptr, &ptr, +/// size*sizeof(Ty), maptype, null); +/// if (res != 0) // Data mapping failed. +/// return res; +/// } +/// } +/// \endcode +llvm::Function *CGOpenMPRuntime::emitUDMapperFunc(const OMPDeclareMapperDecl *D, + bool NoWait) { + ASTContext &C = CGM.getContext(); + QualType Ty = D->getType(); + QualType PtrTy = C.getPointerType(Ty).withRestrict(); + QualType SizeTy = C.getSizeType(); + QualType Int64Ty = C.getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/true); + auto *MapperVarDecl = + cast(cast(D->getMapperVarRef())->getDecl()); + SourceLocation Loc = D->getLocation(); + CharUnits ElementSize = C.getTypeSizeInChars(Ty); + + // Prepare mapper function arguments and attributes. + ImplicitParamDecl DeviceIdArg(C, Int64Ty, ImplicitParamDecl::Other); + ImplicitParamDecl BasePtrArg(C, /*DC=*/nullptr, MapperVarDecl->getLocation(), + /*Id=*/nullptr, C.VoidPtrTy, + ImplicitParamDecl::Other); + ImplicitParamDecl PtrArg(C, /*DC=*/nullptr, MapperVarDecl->getLocation(), + /*Id=*/nullptr, C.VoidPtrTy, + ImplicitParamDecl::Other); + ImplicitParamDecl SizeArg(C, SizeTy, ImplicitParamDecl::Other); + ImplicitParamDecl MapTypeArg(C, Int64Ty, ImplicitParamDecl::Other); + FunctionArgList Args; + Args.push_back(&DeviceIdArg); + Args.push_back(&BasePtrArg); + Args.push_back(&PtrArg); + Args.push_back(&SizeArg); + Args.push_back(&MapTypeArg); + const CGFunctionInfo &FnInfo = + CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.IntTy, Args); + llvm::FunctionType *FnTy = CGM.getTypes().GetFunctionType(FnInfo); + std::string Name = getName( + {"omp_mapper", Ty.getAsString(), D->getName(), NoWait ? "nowait." : ""}); + std::replace(Name.begin(), Name.end(), ' ', '_'); + auto *Fn = llvm::Function::Create(FnTy, llvm::GlobalValue::InternalLinkage, + Name, &CGM.getModule()); + CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, FnInfo); + Fn->removeFnAttr(llvm::Attribute::OptimizeNone); + // Start the mapper function code generation. + CodeGenFunction MapperCGF(CGM); + MapperCGF.StartFunction(GlobalDecl(), C.IntTy, Fn, FnInfo, Args, Loc, Loc); + // Initiate the return value to 0, which represents success. + llvm::Value *SuccessRetVal = llvm::ConstantInt::getNullValue(CGM.IntTy); + MapperCGF.EmitStoreOfScalar(SuccessRetVal, MapperCGF.ReturnValue, + /*Volatile=*/false, C.IntTy); + // Compute the starting and end addreses of array elements. + llvm::Value *Size = MapperCGF.EmitLoadOfScalar( + MapperCGF.GetAddrOfLocalVar(&SizeArg), /*Volatile=*/false, + C.getPointerType(SizeTy), Loc); + llvm::Value *Ptr = MapperCGF.GetAddrOfLocalVar(&PtrArg).getPointer(); + llvm::Value *PtrBegin = MapperCGF.Builder.CreateBitCast( + Ptr, CGM.getTypes().ConvertTypeForMem(C.getPointerType(PtrTy))); + llvm::Value *PtrEnd = MapperCGF.Builder.CreateGEP(PtrBegin, Size); + llvm::Value *NullMapperArrayArg = + llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy); + llvm::Value *MapType = MapperCGF.EmitLoadOfScalar( + MapperCGF.GetAddrOfLocalVar(&MapTypeArg), /*Volatile=*/false, + C.getPointerType(Int64Ty), Loc); + // Prepare some common arguments. + llvm::Value *DeviceID = MapperCGF.EmitLoadOfScalar( + MapperCGF.GetAddrOfLocalVar(&DeviceIdArg), /*Volatile=*/false, + C.getPointerType(Int64Ty), Loc); + llvm::Value *BasePtr = MapperCGF.GetAddrOfLocalVar(&BasePtrArg).getPointer(); + + // Evaluate if this is an array section. + llvm::BasicBlock *IsNotDeleteBB = + MapperCGF.createBasicBlock("omp.arrayinit.evaldelete"); + llvm::BasicBlock *ArrayInitBB = MapperCGF.createBasicBlock("omp.arrayinit"); + llvm::BasicBlock *HeadBB = MapperCGF.createBasicBlock("omp.arraymap.head"); + llvm::Value *IsArray = MapperCGF.Builder.CreateICmpSGE( + Size, MapperCGF.Builder.getIntN(C.getTypeSize(SizeTy), 1), + "omp.arrayinit.isarray"); + MapperCGF.Builder.CreateCondBr(IsArray, IsNotDeleteBB, HeadBB); + // Evaluate if we are going to delete this section. + MapperCGF.EmitBlock(IsNotDeleteBB); + llvm::Value *DeleteBit = MapperCGF.Builder.CreateAnd( + MapType, + MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_DELETE)); + llvm::Value *IsNotDelete = + MapperCGF.Builder.CreateIsNull(DeleteBit, "omp.arrayinit.notdelete"); + MapperCGF.Builder.CreateCondBr(IsNotDelete, ArrayInitBB, HeadBB); + + // Allocate the space if this is an array. + MapperCGF.EmitBlock(ArrayInitBB); + llvm::Value *InitReturn = + emitUDMapperArrayInitOrDel(MapperCGF, DeviceID, BasePtr, Ptr, Size, + MapType, ElementSize, /*IsInit=*/true, NoWait); + + // Jump to the function end if the return value indicates data mapping failed. + llvm::BasicBlock *InitErrorBB = + MapperCGF.createBasicBlock("omp.arrayinit.error"); + llvm::BasicBlock *ExitBB = MapperCGF.createBasicBlock("omp.arraymap.exit"); + llvm::BasicBlock *DoneBB = MapperCGF.createBasicBlock("omp.done"); + llvm::Value *IsInitFailed = MapperCGF.Builder.CreateIsNotNull(InitReturn); + MapperCGF.Builder.CreateCondBr(IsInitFailed, InitErrorBB, HeadBB); + MapperCGF.EmitBlock(InitErrorBB); + MapperCGF.EmitStoreOfScalar(InitReturn, MapperCGF.ReturnValue, + /*Volatile=*/false, C.IntTy); + MapperCGF.Builder.CreateBr(DoneBB); + + // Emit a for loop to iterate through SizeArg of elements and map all of them. + + // Emit the loop header block. + MapperCGF.EmitBlock(HeadBB); + llvm::BasicBlock *BodyBB = MapperCGF.createBasicBlock("omp.arraymap.body"); + llvm::BasicBlock *ErrorBB = MapperCGF.createBasicBlock("omp.arraymap.error"); + llvm::BasicBlock *CorrectBB = + MapperCGF.createBasicBlock("omp.arraymap.correct"); + // Evaluate whether the initial condition is satisfied. + llvm::Value *IsEmpty = + MapperCGF.Builder.CreateICmpEQ(PtrBegin, PtrEnd, "omp.arraymap.isempty"); + MapperCGF.Builder.CreateCondBr(IsEmpty, DoneBB, BodyBB); + llvm::BasicBlock *EntryBB = MapperCGF.Builder.GetInsertBlock(); + + // Emit the loop body block. + MapperCGF.EmitBlock(BodyBB); + llvm::PHINode *PtrPHI = MapperCGF.Builder.CreatePHI( + PtrBegin->getType(), 2, "omp.arraymap.ptrcurrent"); + PtrPHI->addIncoming(PtrBegin, EntryBB); + Address PtrCurrent = + Address(PtrPHI, MapperCGF.GetAddrOfLocalVar(&PtrArg) + .getAlignment() + .alignmentOfArrayElement(ElementSize)); + // Privatize the declared variable of mapper to be the current array element. + CodeGenFunction::OMPPrivateScope Scope(MapperCGF); + Scope.addPrivate(MapperVarDecl, [&MapperCGF, PtrCurrent, PtrTy]() { + return MapperCGF + .EmitLoadOfPointerLValue(PtrCurrent, PtrTy->castAs()) + .getAddress(); + }); + (void)Scope.Privatize(); + + // Get map clause information. + // Fill up the arrays with all the mapped variables. + MappableExprsHandler::MapBaseValuesArrayTy BasePointers; + MappableExprsHandler::MapValuesArrayTy Pointers; + MappableExprsHandler::MapValuesArrayTy Sizes; + MappableExprsHandler::MapFlagsArrayTy MapTypes; + MappableExprsHandler MEHandler(*D, MapperCGF); + MEHandler.generateAllInfoForMapper(BasePointers, Pointers, Sizes, MapTypes); + // Fill up the arrays and create the arguments. + TargetDataInfo Info; + emitOffloadingArrays(MapperCGF, BasePointers, Pointers, Sizes, MapTypes, Info, + MapType); + llvm::Value *BasePointersArrayArg = nullptr; + llvm::Value *PointersArrayArg = nullptr; + llvm::Value *SizesArrayArg = nullptr; + llvm::Value *MapTypesArrayArg = nullptr; + emitOffloadingArraysArgument(MapperCGF, BasePointersArrayArg, + PointersArrayArg, SizesArrayArg, + MapTypesArrayArg, Info); + + // Call the runtime API __tgt_target_data_mapper(_nowait) to map data. + llvm::Value *PointerNum = MapperCGF.Builder.getInt32(Info.NumberOfPtrs); + llvm::Value *OffloadingArgs[] = { + DeviceID, PointerNum, BasePointersArrayArg, PointersArrayArg, + SizesArrayArg, MapTypesArrayArg, NullMapperArrayArg}; + llvm::Value *Return = MapperCGF.EmitRuntimeCall( + createRuntimeFunction(NoWait ? OMPRTL__tgt_target_data_mapper_nowait + : OMPRTL__tgt_target_data_mapper), + OffloadingArgs); + + // Break the loop if the return value indicates data mapping failed. + llvm::Value *IsFailed = MapperCGF.Builder.CreateIsNotNull(Return); + MapperCGF.Builder.CreateCondBr(IsFailed, ErrorBB, CorrectBB); + MapperCGF.EmitBlock(ErrorBB); + MapperCGF.EmitStoreOfScalar(Return, MapperCGF.ReturnValue, + /*Volatile=*/false, C.IntTy); + MapperCGF.Builder.CreateBr(DoneBB); + + // Update the pointer to point to the next element that needs to be mapped, + // and check whether we have mapped all elements. + MapperCGF.EmitBlock(CorrectBB); + llvm::Value *PtrNext = MapperCGF.Builder.CreateConstGEP1_32( + PtrPHI, /*Idx0=*/1, "omp.arraymap.next"); + PtrPHI->addIncoming(PtrNext, CorrectBB); + llvm::Value *IsDone = + MapperCGF.Builder.CreateICmpEQ(PtrNext, PtrEnd, "omp.arraymap.isdone"); + MapperCGF.Builder.CreateCondBr(IsDone, ExitBB, BodyBB); + + // Delete the array section if specified by the input map type. + MapperCGF.EmitBlock(ExitBB); + llvm::BasicBlock *IsDeleteBB = + MapperCGF.createBasicBlock("omp.arraydel.evaldelete"); + llvm::BasicBlock *ArrayDelBB = MapperCGF.createBasicBlock("omp.arraydel"); + MapperCGF.Builder.CreateCondBr(IsArray, IsDeleteBB, DoneBB); + MapperCGF.EmitBlock(IsDeleteBB); + llvm::Value *DB = MapperCGF.Builder.CreateAnd( + MapType, + MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_DELETE)); + llvm::Value *IsDelete = + MapperCGF.Builder.CreateIsNotNull(DB, "omp.arraydel.isdelete"); + MapperCGF.Builder.CreateCondBr(IsDelete, ArrayDelBB, DoneBB); + MapperCGF.EmitBlock(ArrayDelBB); + llvm::Value *DelReturn = emitUDMapperArrayInitOrDel( + MapperCGF, DeviceID, BasePtr, Ptr, Size, MapType, ElementSize, + /*IsInit=*/false, NoWait); + + // Jump to the function end if the return value indicates data mapping failed. + llvm::BasicBlock *DelErrorBB = + MapperCGF.createBasicBlock("omp.arraydel.error"); + llvm::Value *IsDelFailed = MapperCGF.Builder.CreateIsNotNull(DelReturn); + MapperCGF.Builder.CreateCondBr(IsDelFailed, DelErrorBB, DoneBB); + MapperCGF.EmitBlock(DelErrorBB); + MapperCGF.EmitStoreOfScalar(DelReturn, MapperCGF.ReturnValue, + /*Volatile=*/false, C.IntTy); + + // Emit the function exit block. + MapperCGF.EmitBlock(DoneBB, /*IsFinished=*/true); + MapperCGF.FinishFunction(); + return Fn; +} + +// Emit the array initialization or deletion portion for user-defined mapper +// code generation. +llvm::Value *CGOpenMPRuntime::emitUDMapperArrayInitOrDel( + CodeGenFunction &MapperCGF, llvm::Value *DeviceID, llvm::Value *BasePtr, + llvm::Value *Ptr, llvm::Value *Size, llvm::Value *MapType, + CharUnits ElementSize, bool IsInit, bool NoWait) { + ASTContext &C = CGM.getContext(); + QualType Int64Ty = C.getIntTypeForBitwidth(/*DestWidth*/ 64, /*Signed*/ true); + std::string Prefix = IsInit ? ".init" : ".del"; + // Prepare the size argument. + unsigned SizeTyWidth = C.getTypeSize(C.getSizeType()); + llvm::Value *ArraySize = MapperCGF.Builder.CreateMul( + Size, MapperCGF.Builder.getIntN(SizeTyWidth, ElementSize.getQuantity())); + llvm::APInt PointerNumAP(32, 1, /*isSigned=*/true); + QualType SizeArrayType = + C.getConstantArrayType(C.getSizeType(), PointerNumAP, ArrayType::Normal, + /*IndexTypeQuals=*/0); + llvm::Value *SizesArrayStorage = + MapperCGF.CreateMemTemp(SizeArrayType, Prefix + ".offload_sizes") + .getPointer(); + llvm::Value *SizesArrayArg = MapperCGF.Builder.CreateConstInBoundsGEP2_32( + llvm::ArrayType::get(CGM.SizeTy, 1), SizesArrayStorage, /*Idx0=*/0, + /*Idx1=*/0); + Address SizesArrayAddr(SizesArrayArg, C.getTypeAlignInChars(C.getSizeType())); + MapperCGF.EmitStoreOfScalar(ArraySize, SizesArrayAddr, /*Volatile=*/false, + C.getSizeType()); + // Prepare the map type argument. + QualType MapArrayType = + C.getConstantArrayType(Int64Ty, PointerNumAP, ArrayType::Normal, + /*IndexTypeQuals=*/0); + llvm::Value *MapTypeArrayStorage = + MapperCGF.CreateMemTemp(MapArrayType, Prefix + ".offload_maptypes") + .getPointer(); + llvm::Value *MapTypeArg = MapperCGF.Builder.CreateConstInBoundsGEP2_32( + llvm::ArrayType::get(CGM.Int64Ty, 1), MapTypeArrayStorage, /*Idx0=*/0, + /*Idx1=*/0); + Address MapTypeArrayAddr(MapTypeArg, C.getTypeAlignInChars(Int64Ty)); + // Remove OMP_MAP_TO and OMP_MAP_FROM from the map type, so that it achieves + // memory allocation/deletion purpose only. + llvm::Value *ArrayMapType = MapperCGF.Builder.CreateAnd( + MapType, + MapperCGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO | + MappableExprsHandler::OMP_MAP_FROM))); + MapperCGF.EmitStoreOfScalar(ArrayMapType, MapTypeArrayAddr, + /*Volatile=*/false, Int64Ty); + llvm::Value *NullMapperArrayArg = + llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy); + llvm::Value *OffloadingArgs[] = { + DeviceID, + /*arg_num*/ MapperCGF.Builder.getInt32(1), + BasePtr, + Ptr, + SizesArrayArg, + MapTypeArg, + NullMapperArrayArg}; + return MapperCGF.EmitRuntimeCall( + createRuntimeFunction(NoWait ? OMPRTL__tgt_target_data_mapper_nowait + : OMPRTL__tgt_target_data_mapper), + OffloadingArgs); +} + void CGOpenMPRuntime::emitTargetNumIterationsCall( CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *Device, const llvm::function_ref(Member)) { if (Ctx->DeclMustBeEmitted(DRD)) Builder->EmitGlobal(DRD); + } else if (auto *DMD = dyn_cast(Member)) { + if (Ctx->DeclMustBeEmitted(DMD)) + Builder->EmitGlobal(DMD); } } } Index: test/OpenMP/declare_mapper_codegen.cpp =================================================================== --- test/OpenMP/declare_mapper_codegen.cpp +++ test/OpenMP/declare_mapper_codegen.cpp @@ -1,92 +1,770 @@ -///==========================================================================/// -// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s -// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s -// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s -// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s -// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s - -// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s -// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s -// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s -// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s -// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s -// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY0 %s - // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} // expected-no-diagnostics #ifndef HEADER #define HEADER +///==========================================================================/// +// RUN: %clang_cc1 -DCK0 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK0 --check-prefix CK0-64 %s +// RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK0 --check-prefix CK0-64 %s +// RUN: %clang_cc1 -DCK0 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK0 --check-prefix CK0-32 %s +// RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK0 --check-prefix CK0-32 %s + +// RUN: %clang_cc1 -DCK0 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK0 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s + +#ifdef CK0 + +// CK0-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0 +// CK0-64: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i64 16] +// CK0-32: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i32 8] +// CK0: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35] +// CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 16] +// CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 8] +// CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33] +// CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 16] +// CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 8] +// CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34] + class C { public: int a; + double *b; }; -#pragma omp declare mapper(id: C s) map(s.a) +#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2]) + +// Synchronous version of mapper function. +// CK0-LABEL: define {{.*}}i32 @.omp_mapper.class_C.id{{.*}}(i64, i8*, i8*, i{{64|32}}, i64) +// CK0-DAG: store i64 %0, i64* [[DIDADDR:%[^,]+]] +// CK0-DAG: store i[[sz:64|32]] %3, i{{64|32}}* [[SIZEADDR:%[^,]+]] +// CK0-DAG: store i64 %4, i64* [[TYPEADDR:%[^,]+]] +// CK0-DAG: store i8* %1, i8** [[BPTRADDR:%[^,]+]] +// CK0-DAG: store i8* %2, i8** [[VPTRADDR:%[^,]+]] +// CK0-DAG: store i32 0, i32* %retval +// CK0-DAG: [[SIZE:%.+]] = load i[[sz]], i[[sz]]* [[SIZEADDR]] +// CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] +// CK0-DAG: [[DID:%.+]] = load i64, i64* [[DIDADDR]] +// CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** +// CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i[[sz]] [[SIZE]] +// CK0: [[ISARRAY:%.+]] = icmp sge i[[sz]] [[SIZE]], 1 +// CK0: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] + +// CK0: [[INITEVALDEL]] +// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK0: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK0: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] +// CK0: [[INIT]] +// CK0-64: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 16 +// CK0-32: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 8 +// CK0-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[ISIZEADDR:[^,]+]] +// CK0-DAG: [[ISIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[ISIZE:%.+]], i32 0, i32 0 +// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK0-DAG: store i64 [[ITYPE]], i64* [[ITYPEADDR:[^,]+]] +// CK0-DAG: [[ITYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[ITYPE:%.+]], i32 0, i32 0 +// CK0: [[IRES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[ISIZEADDR]], i64* [[ITYPEADDR]], i8** null) +// CK0: [[ISINITERR:%.+]] = icmp ne i32 [[IRES]], 0 +// CK0: br i1 [[ISINITERR]], label %[[INITERR:[^,]+]], label %[[LHEAD:[^,]+]] +// CK0: [[INITERR]] +// CK0: store i32 [[IRES]], i32* %retval +// CK0: br label %[[DONE:[^,]+]] + +// CK0: [[LHEAD]] +// CK0: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK0: br i1 [[ISEMPTY]], label %[[DONE]], label %[[LBODY:[^,]+]] +// CK0: [[LBODY]] +// CK0: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK0: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] +// CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0 +// CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK0-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]] +// CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz]] 0 +// CK0-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1 +// CK0-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8* +// CK0-DAG: [[BENDV:%.+]] = bitcast double** [[BEND]] to i8* +// CK0-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64 +// CK0-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64 +// CK0-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]] +// CK0-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK0-32-DAG: [[CUSIZE32:%.+]] = trunc i64 [[CUSIZE]] to i32 +// CK0-64-DAG: store i[[sz]] [[CUSIZE]], i[[sz]]* [[SIZEADDR0:%[^,]+]] +// CK0-32-DAG: store i[[sz]] [[CUSIZE32]], i[[sz]]* [[SIZEADDR0:%[^,]+]] +// CK0-DAG: [[SIZEADDR0]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR:%[^,]+]], i32 0, i32 0 +// CK0-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR0BC:%[^,]+]] +// CK0-DAG: [[BPTRADDR0BC]] = bitcast i8** [[BPTRADDR0:%.+]] to %class.C** +// CK0-DAG: [[BPTRADDR0]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR:%[^,]+]], i32 0, i32 0 +// CK0-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR0BC:%[^,]+]] +// CK0-DAG: [[PTRADDR0BC]] = bitcast i8** [[PTRADDR0:%.+]] to i32** +// CK0-DAG: [[PTRADDR0]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR:%[^,]+]], i32 0, i32 0 +// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK0-DAG: [[ALLOC]] +// CK0-DAG: br label %[[TYEND:[^,]+]] +// CK0-DAG: [[ALLOCELSE]] +// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK0-DAG: [[TO]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TOELSE]] +// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK0-DAG: [[FROM]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TYEND]] +// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ 32, %[[ALLOC]] ], [ 32, %[[TO]] ], [ 32, %[[FROM]] ], [ 32, %[[TOELSE]] ] +// CK0-DAG: store i64 [[PHITYPE0]], i64* [[TYPEADDR0:%[^,]+]] +// CK0-DAG: [[TYPEADDR0]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR:%[^,]+]], i32 0, i32 0 +// CK0-DAG: [[BPTRADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 1 +// CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast i8** [[BPTRADDR1]] to %class.C** +// CK0-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR1BC]] +// CK0-DAG: [[PTRADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 1 +// CK0-DAG: [[PTRADDR1BC:%.+]] = bitcast i8** [[PTRADDR1]] to i32** +// CK0-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR1BC]] +// CK0-DAG: [[SIZEADDR1:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 1 +// CK0-DAG: store i[[sz]] 4, i[[sz]]* [[SIZEADDR1]] +// CK0-DAG: [[TYPEADDR1:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 1 +// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK0-DAG: [[ALLOC]] +// CK0-DAG: br label %[[TYEND:[^,]+]] +// CK0-DAG: [[ALLOCELSE]] +// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK0-DAG: [[TO]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TOELSE]] +// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK0-DAG: [[FROM]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TYEND]] +// 0x1,000,000,000; 0x1,000,000,001; 0x1,000,000,002; 0x1,000,000,003 +// CK0-DAG: [[TYPE1:%.+]] = phi i64 [ 281474976710656, %[[ALLOC]] ], [ 281474976710657, %[[TO]] ], [ 281474976710658, %[[FROM]] ], [ 281474976710659, %[[TOELSE]] ] +// CK0-DAG: store i64 [[TYPE1]], i64* [[TYPEADDR1]] +// CK0-DAG: [[BPTRADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 2 +// CK0-DAG: [[BPTRADDR2BC:%.+]] = bitcast i8** [[BPTRADDR2]] to double*** +// CK0-DAG: store double** [[BBEGIN]], double*** [[BPTRADDR2BC]] +// CK0-DAG: [[PTRADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 2 +// CK0-DAG: [[PTRADDR2BC:%.+]] = bitcast i8** [[PTRADDR2]] to double** +// CK0-DAG: store double* [[BARRBEGINGEP]], double** [[PTRADDR2BC]] +// CK0-DAG: [[SIZEADDR2:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 2 +// CK0-DAG: store i[[sz]] 16, i[[sz]]* [[SIZEADDR2]] +// CK0-DAG: [[TYPEADDR2:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 2 +// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK0-DAG: [[ALLOC]] +// CK0-DAG: br label %[[TYEND:[^,]+]] +// CK0-DAG: [[ALLOCELSE]] +// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK0-DAG: [[TO]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TOELSE]] +// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK0-DAG: [[FROM]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TYEND]] +// 0x1,000,000,010; 0x1,000,000,011; 0x1,000,000,012; 0x1,000,000,013 +// CK0-DAG: [[TYPE2:%.+]] = phi i64 [ 281474976710672, %[[ALLOC]] ], [ 281474976710673, %[[TO]] ], [ 281474976710674, %[[FROM]] ], [ 281474976710675, %[[TOELSE]] ] +// CK0-DAG: store i64 [[TYPE2]], i64* [[TYPEADDR2]] +// CK0-DAG: [[ARGBPTR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 0 +// CK0-DAG: [[ARGPTR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 0 +// CK0-DAG: [[ARGSIZE:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 0 +// CK0-DAG: [[ARGTYPE:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 0 +// CK0: [[RES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 3, i8** [[ARGBPTR]], i8** [[ARGPTR]], i[[sz]]* [[ARGSIZE]], i64* [[ARGTYPE]], i8** null) +// CK0: [[ISERR:%.+]] = icmp ne i32 [[RES]], 0 +// CK0: br i1 [[ISERR]], label %[[LERR:[^,]+]], label %[[LCORRECT]] +// CK0: [[LERR]] +// CK0: store i32 [[RES]], i32* %retval +// CK0: br label %[[DONE]] +// CK0: [[LCORRECT]] +// CK0: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 +// CK0: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] -// CHECK-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l54.region_id = weak constant i8 0 +// CK0: [[LEXIT]] +// CK0: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] +// CK0: [[EVALDEL]] +// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK0: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK0: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] +// CK0: [[DEL]] +// CK0-64: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 16 +// CK0-32: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 8 +// CK0-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[DSIZEADDR:%[^,]+]] +// CK0-DAG: [[DSIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[DSIZE:%.+]], i32 0, i32 0 +// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK0-DAG: store i64 [[DTYPE]], i64* [[DTYPEADDR:%[^,]+]] +// CK0-DAG: [[DTYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[DTYPE:%.+]], i32 0, i32 0 +// CK0: [[DRES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[DSIZEADDR]], i64* [[DTYPEADDR]], i8** null) +// CK0: [[ISDELERR:%.+]] = icmp ne i32 [[DRES]], 0 +// CK0: br i1 [[ISDELERR]], label %[[DELERR:[^,]+]], label %[[DONE]] +// CK0: [[DELERR]] +// CK0: store i32 [[DRES]], i32* %retval +// CK0: br label %[[DONE]] +// CK0: [[DONE]] +// CK0: [[RET:%.+]] = load i32, i32* %retval +// CK0: ret i32 [[RET]] -// CHECK: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] -// CHECK: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35] -// CHECK: [[TSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4] -// CHECK: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33] -// CHECK: [[FSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4] -// CHECK: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34] -// CHECK-LABEL: foo{{.*}}( +// Asynchronous version of mapper function. +// CK0-LABEL: define {{.*}}i32 @.omp_mapper.class_C.id{{.*}}nowait{{.*}}(i64, i8*, i8*, i{{64|32}}, i64) +// CK0-DAG: store i64 %0, i64* [[DIDADDR:%[^,]+]] +// CK0-DAG: store i[[sz:64|32]] %3, i{{64|32}}* [[SIZEADDR:%[^,]+]] +// CK0-DAG: store i64 %4, i64* [[TYPEADDR:%[^,]+]] +// CK0-DAG: store i8* %1, i8** [[BPTRADDR:%[^,]+]] +// CK0-DAG: store i8* %2, i8** [[VPTRADDR:%[^,]+]] +// CK0-DAG: store i32 0, i32* %retval +// CK0-DAG: [[SIZE:%.+]] = load i[[sz]], i[[sz]]* [[SIZEADDR]] +// CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] +// CK0-DAG: [[DID:%.+]] = load i64, i64* [[DIDADDR]] +// CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** +// CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i[[sz]] [[SIZE]] +// CK0: [[ISARRAY:%.+]] = icmp sge i[[sz]] [[SIZE]], 1 +// CK0: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] + +// CK0: [[INITEVALDEL]] +// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK0: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK0: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] +// CK0: [[INIT]] +// CK0-64: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 16 +// CK0-32: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 8 +// CK0-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[ISIZEADDR:[^,]+]] +// CK0-DAG: [[ISIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[ISIZE:%.+]], i32 0, i32 0 +// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK0-DAG: store i64 [[ITYPE]], i64* [[ITYPEADDR:[^,]+]] +// CK0-DAG: [[ITYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[ITYPE:%.+]], i32 0, i32 0 +// CK0: [[IRES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[ISIZEADDR]], i64* [[ITYPEADDR]], i8** null) +// CK0: [[ISINITERR:%.+]] = icmp ne i32 [[IRES]], 0 +// CK0: br i1 [[ISINITERR]], label %[[INITERR:[^,]+]], label %[[LHEAD:[^,]+]] +// CK0: [[INITERR]] +// CK0: store i32 [[IRES]], i32* %retval +// CK0: br label %[[DONE:[^,]+]] + +// CK0: [[LHEAD]] +// CK0: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK0: br i1 [[ISEMPTY]], label %[[DONE]], label %[[LBODY:[^,]+]] +// CK0: [[LBODY]] +// CK0: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK0: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] +// CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0 +// CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK0-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]] +// CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz]] 0 +// CK0-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1 +// CK0-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8* +// CK0-DAG: [[BENDV:%.+]] = bitcast double** [[BEND]] to i8* +// CK0-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64 +// CK0-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64 +// CK0-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]] +// CK0-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK0-32-DAG: [[CUSIZE32:%.+]] = trunc i64 [[CUSIZE]] to i32 +// CK0-64-DAG: store i[[sz]] [[CUSIZE]], i[[sz]]* [[SIZEADDR0:%[^,]+]] +// CK0-32-DAG: store i[[sz]] [[CUSIZE32]], i[[sz]]* [[SIZEADDR0:%[^,]+]] +// CK0-DAG: [[SIZEADDR0]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR:%[^,]+]], i32 0, i32 0 +// CK0-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR0BC:%[^,]+]] +// CK0-DAG: [[BPTRADDR0BC]] = bitcast i8** [[BPTRADDR0:%.+]] to %class.C** +// CK0-DAG: [[BPTRADDR0]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR:%[^,]+]], i32 0, i32 0 +// CK0-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR0BC:%[^,]+]] +// CK0-DAG: [[PTRADDR0BC]] = bitcast i8** [[PTRADDR0:%.+]] to i32** +// CK0-DAG: [[PTRADDR0]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR:%[^,]+]], i32 0, i32 0 +// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK0-DAG: [[ALLOC]] +// CK0-DAG: br label %[[TYEND:[^,]+]] +// CK0-DAG: [[ALLOCELSE]] +// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK0-DAG: [[TO]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TOELSE]] +// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK0-DAG: [[FROM]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TYEND]] +// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ 32, %[[ALLOC]] ], [ 32, %[[TO]] ], [ 32, %[[FROM]] ], [ 32, %[[TOELSE]] ] +// CK0-DAG: store i64 [[PHITYPE0]], i64* [[TYPEADDR0:%[^,]+]] +// CK0-DAG: [[TYPEADDR0]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR:%[^,]+]], i32 0, i32 0 +// CK0-DAG: [[BPTRADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 1 +// CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast i8** [[BPTRADDR1]] to %class.C** +// CK0-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR1BC]] +// CK0-DAG: [[PTRADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 1 +// CK0-DAG: [[PTRADDR1BC:%.+]] = bitcast i8** [[PTRADDR1]] to i32** +// CK0-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR1BC]] +// CK0-DAG: [[SIZEADDR1:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 1 +// CK0-DAG: store i[[sz]] 4, i[[sz]]* [[SIZEADDR1]] +// CK0-DAG: [[TYPEADDR1:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 1 +// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK0-DAG: [[ALLOC]] +// CK0-DAG: br label %[[TYEND:[^,]+]] +// CK0-DAG: [[ALLOCELSE]] +// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK0-DAG: [[TO]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TOELSE]] +// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK0-DAG: [[FROM]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TYEND]] +// 0x1,000,000,000; 0x1,000,000,001; 0x1,000,000,002; 0x1,000,000,003 +// CK0-DAG: [[TYPE1:%.+]] = phi i64 [ 281474976710656, %[[ALLOC]] ], [ 281474976710657, %[[TO]] ], [ 281474976710658, %[[FROM]] ], [ 281474976710659, %[[TOELSE]] ] +// CK0-DAG: store i64 [[TYPE1]], i64* [[TYPEADDR1]] +// CK0-DAG: [[BPTRADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 2 +// CK0-DAG: [[BPTRADDR2BC:%.+]] = bitcast i8** [[BPTRADDR2]] to double*** +// CK0-DAG: store double** [[BBEGIN]], double*** [[BPTRADDR2BC]] +// CK0-DAG: [[PTRADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 2 +// CK0-DAG: [[PTRADDR2BC:%.+]] = bitcast i8** [[PTRADDR2]] to double** +// CK0-DAG: store double* [[BARRBEGINGEP]], double** [[PTRADDR2BC]] +// CK0-DAG: [[SIZEADDR2:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 2 +// CK0-DAG: store i[[sz]] 16, i[[sz]]* [[SIZEADDR2]] +// CK0-DAG: [[TYPEADDR2:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 2 +// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK0-DAG: [[ALLOC]] +// CK0-DAG: br label %[[TYEND:[^,]+]] +// CK0-DAG: [[ALLOCELSE]] +// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK0-DAG: [[TO]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TOELSE]] +// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK0-DAG: [[FROM]] +// CK0-DAG: br label %[[TYEND]] +// CK0-DAG: [[TYEND]] +// 0x1,000,000,010; 0x1,000,000,011; 0x1,000,000,012; 0x1,000,000,013 +// CK0-DAG: [[TYPE2:%.+]] = phi i64 [ 281474976710672, %[[ALLOC]] ], [ 281474976710673, %[[TO]] ], [ 281474976710674, %[[FROM]] ], [ 281474976710675, %[[TOELSE]] ] +// CK0-DAG: store i64 [[TYPE2]], i64* [[TYPEADDR2]] +// CK0-DAG: [[ARGBPTR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[LBPTRADDR]], i32 0, i32 0 +// CK0-DAG: [[ARGPTR:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRADDR]], i32 0, i32 0 +// CK0-DAG: [[ARGSIZE:%.+]] = getelementptr inbounds [3 x i[[sz]]], [3 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 0 +// CK0-DAG: [[ARGTYPE:%.+]] = getelementptr inbounds [3 x i64], [3 x i64]* [[TYPEADDR]], i32 0, i32 0 +// CK0: [[RES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 3, i8** [[ARGBPTR]], i8** [[ARGPTR]], i[[sz]]* [[ARGSIZE]], i64* [[ARGTYPE]], i8** null) +// CK0: [[ISERR:%.+]] = icmp ne i32 [[RES]], 0 +// CK0: br i1 [[ISERR]], label %[[LERR:[^,]+]], label %[[LCORRECT]] +// CK0: [[LERR]] +// CK0: store i32 [[RES]], i32* %retval +// CK0: br label %[[DONE]] +// CK0: [[LCORRECT]] +// CK0: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 +// CK0: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] + +// CK0: [[LEXIT]] +// CK0: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] +// CK0: [[EVALDEL]] +// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK0: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK0: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] +// CK0: [[DEL]] +// CK0-64: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 16 +// CK0-32: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 8 +// CK0-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[DSIZEADDR:%[^,]+]] +// CK0-DAG: [[DSIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[DSIZE:%.+]], i32 0, i32 0 +// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK0-DAG: store i64 [[DTYPE]], i64* [[DTYPEADDR:%[^,]+]] +// CK0-DAG: [[DTYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[DTYPE:%.+]], i32 0, i32 0 +// CK0: [[DRES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[DSIZEADDR]], i64* [[DTYPEADDR]], i8** null) +// CK0: [[ISDELERR:%.+]] = icmp ne i32 [[DRES]], 0 +// CK0: br i1 [[ISDELERR]], label %[[DELERR:[^,]+]], label %[[DONE]] +// CK0: [[DELERR]] +// CK0: store i32 [[DRES]], i32* %retval +// CK0: br label %[[DONE]] +// CK0: [[DONE]] +// CK0: [[RET:%.+]] = load i32, i32* %retval +// CK0: ret i32 [[RET]] + + +// CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}} void foo(int a){ int i = a; C c; c.a = a; - // CHECK-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) - // CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 - // CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 - // CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 - // CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 - // CHECK-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C** - // CHECK-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C** - // CHECK-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]] - // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[CP1]] - // CHECK: call void [[KERNEL:@.+]](%class.C* [[VAL]]) + // CK0-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) + // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 + // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 + // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 + // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 + // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C** + // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C** + // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]] + // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]] + // CK0: call void [[KERNEL:@.+]](%class.C* [[VAL]]) #pragma omp target map(mapper(id),tofrom: c) { ++c.a; } - // CHECK-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}) - // CHECK-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 - // CHECK-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 - // CHECK-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0 - // CHECK-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0 - // CHECK-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C** - // CHECK-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C** - // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]] - // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]] + // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}) + // CK0-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C** + // CK0-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C** + // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]] + // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]] #pragma omp target update to(mapper(id): c) - // CHECK-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}) - // CHECK-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 - // CHECK-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 - // CHECK-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0 - // CHECK-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0 - // CHECK-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C** - // CHECK-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C** - // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]] - // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]] + // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}) + // CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C** + // CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C** + // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]] + // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]] #pragma omp target update from(mapper(id): c) } -// CHECK: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]]) -// CHECK: [[ADDR:%.+]] = alloca %class.C*, -// CHECK: store %class.C* [[ARG]], %class.C** [[ADDR]] -// CHECK: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]] -// CHECK: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0 -// CHECK: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]] -// CHECK: {{.+}} = add nsw i32 [[VAL]], 1 -// CHECK: } +// CK0: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]]) +// CK0: [[ADDR:%.+]] = alloca %class.C*, +// CK0: store %class.C* [[ARG]], %class.C** [[ADDR]] +// CK0: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]] +// CK0: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0 +// CK0: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]] +// CK0: {{.+}} = add nsw i32 [[VAL]], 1 +// CK0: } + +#endif + + +///==========================================================================/// +// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s +// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s +// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s +// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s + +// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s + +#ifdef CK1 + +template +class C { +public: + T a; +}; + +#pragma omp declare mapper(id: C s) map(s.a) + +// Synchronous version of mapper function. +// CK1-LABEL: define {{.*}}i32 @".omp_mapper.C.id{{.*}}(i64, i8*, i8*, i{{64|32}}, i64) +// CK1-DAG: store i64 %0, i64* [[DIDADDR:%[^,]+]] +// CK1-DAG: store i[[sz:64|32]] %3, i{{64|32}}* [[SIZEADDR:%[^,]+]] +// CK1-DAG: store i64 %4, i64* [[TYPEADDR:%[^,]+]] +// CK1-DAG: store i8* %1, i8** [[BPTRADDR:%[^,]+]] +// CK1-DAG: store i8* %2, i8** [[VPTRADDR:%[^,]+]] +// CK1-DAG: store i32 0, i32* %retval +// CK1-DAG: [[SIZE:%.+]] = load i[[sz]], i[[sz]]* [[SIZEADDR]] +// CK1-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] +// CK1-DAG: [[DID:%.+]] = load i64, i64* [[DIDADDR]] +// CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** +// CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i[[sz]] [[SIZE]] +// CK1: [[ISARRAY:%.+]] = icmp sge i[[sz]] [[SIZE]], 1 +// CK1: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] + +// CK1: [[INITEVALDEL]] +// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK1: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK1: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] +// CK1: [[INIT]] +// CK1: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 4 +// CK1-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[ISIZEADDR:[^,]+]] +// CK1-DAG: [[ISIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[ISIZE:%.+]], i32 0, i32 0 +// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK1-DAG: store i64 [[ITYPE]], i64* [[ITYPEADDR:[^,]+]] +// CK1-DAG: [[ITYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[ITYPE:%.+]], i32 0, i32 0 +// CK1: [[IRES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[ISIZEADDR]], i64* [[ITYPEADDR]], i8** null) +// CK1: [[ISINITERR:%.+]] = icmp ne i32 [[IRES]], 0 +// CK1: br i1 [[ISINITERR]], label %[[INITERR:[^,]+]], label %[[LHEAD:[^,]+]] +// CK1: [[INITERR]] +// CK1: store i32 [[IRES]], i32* %retval +// CK1: br label %[[DONE:[^,]+]] + +// CK1: [[LHEAD]] +// CK1: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK1: br i1 [[ISEMPTY]], label %[[DONE]], label %[[LBODY:[^,]+]] +// CK1: [[LBODY]] +// CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK1: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] +// CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0 +// CK1-DAG: [[AEND:%.+]] = getelementptr i32, i32* [[ABEGIN]], i32 1 +// CK1-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8* +// CK1-DAG: [[AENDV:%.+]] = bitcast i32* [[AEND]] to i8* +// CK1-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64 +// CK1-DAG: [[AENDI:%.+]] = ptrtoint i8* [[AENDV]] to i64 +// CK1-DAG: [[CSIZE:%.+]] = sub i64 [[AENDI]], [[ABEGINI]] +// CK1-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK1-32-DAG: [[CUSIZE32:%.+]] = trunc i64 [[CUSIZE]] to i32 +// CK1-64-DAG: store i[[sz]] [[CUSIZE]], i[[sz]]* [[SIZEADDR0:%[^,]+]] +// CK1-32-DAG: store i[[sz]] [[CUSIZE32]], i[[sz]]* [[SIZEADDR0:%[^,]+]] +// CK1-DAG: [[SIZEADDR0]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR:%[^,]+]], i32 0, i32 0 +// CK1-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR0BC:%[^,]+]] +// CK1-DAG: [[BPTRADDR0BC]] = bitcast i8** [[BPTRADDR0:%.+]] to %class.C** +// CK1-DAG: [[BPTRADDR0]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR:%[^,]+]], i32 0, i32 0 +// CK1-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR0BC:%[^,]+]] +// CK1-DAG: [[PTRADDR0BC]] = bitcast i8** [[PTRADDR0:%.+]] to i32** +// CK1-DAG: [[PTRADDR0]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR:%[^,]+]], i32 0, i32 0 +// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK1-DAG: [[ALLOC]] +// CK1-DAG: br label %[[TYEND:[^,]+]] +// CK1-DAG: [[ALLOCELSE]] +// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK1-DAG: [[TO]] +// CK1-DAG: br label %[[TYEND]] +// CK1-DAG: [[TOELSE]] +// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK1-DAG: [[FROM]] +// CK1-DAG: br label %[[TYEND]] +// CK1-DAG: [[TYEND]] +// CK1-DAG: [[PHITYPE0:%.+]] = phi i64 [ 32, %[[ALLOC]] ], [ 32, %[[TO]] ], [ 32, %[[FROM]] ], [ 32, %[[TOELSE]] ] +// CK1-DAG: store i64 [[PHITYPE0]], i64* [[TYPEADDR0:%[^,]+]] +// CK1-DAG: [[TYPEADDR0]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR:%[^,]+]], i32 0, i32 0 +// CK1-DAG: [[BPTRADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR]], i32 0, i32 1 +// CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast i8** [[BPTRADDR1]] to %class.C** +// CK1-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR1BC]] +// CK1-DAG: [[PTRADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR]], i32 0, i32 1 +// CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i8** [[PTRADDR1]] to i32** +// CK1-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR1BC]] +// CK1-DAG: [[SIZEADDR1:%.+]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 1 +// CK1-DAG: store i[[sz]] 4, i[[sz]]* [[SIZEADDR1]] +// CK1-DAG: [[TYPEADDR1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR]], i32 0, i32 1 +// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK1-DAG: [[ALLOC]] +// CK1-DAG: br label %[[TYEND:[^,]+]] +// CK1-DAG: [[ALLOCELSE]] +// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK1-DAG: [[TO]] +// CK1-DAG: br label %[[TYEND]] +// CK1-DAG: [[TOELSE]] +// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK1-DAG: [[FROM]] +// CK1-DAG: br label %[[TYEND]] +// CK1-DAG: [[TYEND]] +// 0x1,000,000,000; 0x1,000,000,001; 0x1,000,000,002; 0x1,000,000,003 +// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ 281474976710656, %[[ALLOC]] ], [ 281474976710657, %[[TO]] ], [ 281474976710658, %[[FROM]] ], [ 281474976710659, %[[TOELSE]] ] +// CK1-DAG: store i64 [[TYPE1]], i64* [[TYPEADDR1]] +// CK1-DAG: [[ARGBPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR]], i32 0, i32 0 +// CK1-DAG: [[ARGPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR]], i32 0, i32 0 +// CK1-DAG: [[ARGSIZE:%.+]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 0 +// CK1-DAG: [[ARGTYPE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR]], i32 0, i32 0 +// CK1: [[RES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 2, i8** [[ARGBPTR]], i8** [[ARGPTR]], i[[sz]]* [[ARGSIZE]], i64* [[ARGTYPE]], i8** null) +// CK1: [[ISERR:%.+]] = icmp ne i32 [[RES]], 0 +// CK1: br i1 [[ISERR]], label %[[LERR:[^,]+]], label %[[LCORRECT]] +// CK1: [[LERR]] +// CK1: store i32 [[RES]], i32* %retval +// CK1: br label %[[DONE]] +// CK1: [[LCORRECT]] +// CK1: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 +// CK1: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] + +// CK1: [[LEXIT]] +// CK1: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] +// CK1: [[EVALDEL]] +// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK1: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK1: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] +// CK1: [[DEL]] +// CK1: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 4 +// CK1-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[DSIZEADDR:%[^,]+]] +// CK1-DAG: [[DSIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[DSIZE:%.+]], i32 0, i32 0 +// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK1-DAG: store i64 [[DTYPE]], i64* [[DTYPEADDR:%[^,]+]] +// CK1-DAG: [[DTYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[DTYPE:%.+]], i32 0, i32 0 +// CK1: [[DRES:%.+]] = call i32 @__tgt_target_data_mapper(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[DSIZEADDR]], i64* [[DTYPEADDR]], i8** null) +// CK1: [[ISDELERR:%.+]] = icmp ne i32 [[DRES]], 0 +// CK1: br i1 [[ISDELERR]], label %[[DELERR:[^,]+]], label %[[DONE]] +// CK1: [[DELERR]] +// CK1: store i32 [[DRES]], i32* %retval +// CK1: br label %[[DONE]] +// CK1: [[DONE]] +// CK1: [[RET:%.+]] = load i32, i32* %retval +// CK1: ret i32 [[RET]] + + +// Asynchronous version of mapper function. +// CK1-LABEL: define {{.*}}i32 @".omp_mapper.C.id{{.*}}nowait{{.*}}(i64, i8*, i8*, i{{64|32}}, i64) +// CK1-DAG: store i64 %0, i64* [[DIDADDR:%[^,]+]] +// CK1-DAG: store i[[sz:64|32]] %3, i{{64|32}}* [[SIZEADDR:%[^,]+]] +// CK1-DAG: store i64 %4, i64* [[TYPEADDR:%[^,]+]] +// CK1-DAG: store i8* %1, i8** [[BPTRADDR:%[^,]+]] +// CK1-DAG: store i8* %2, i8** [[VPTRADDR:%[^,]+]] +// CK1-DAG: store i32 0, i32* %retval +// CK1-DAG: [[SIZE:%.+]] = load i[[sz]], i[[sz]]* [[SIZEADDR]] +// CK1-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] +// CK1-DAG: [[DID:%.+]] = load i64, i64* [[DIDADDR]] +// CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** +// CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i[[sz]] [[SIZE]] +// CK1: [[ISARRAY:%.+]] = icmp sge i[[sz]] [[SIZE]], 1 +// CK1: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] + +// CK1: [[INITEVALDEL]] +// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK1: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK1: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] +// CK1: [[INIT]] +// CK1: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 4 +// CK1-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[ISIZEADDR:[^,]+]] +// CK1-DAG: [[ISIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[ISIZE:%.+]], i32 0, i32 0 +// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK1-DAG: store i64 [[ITYPE]], i64* [[ITYPEADDR:[^,]+]] +// CK1-DAG: [[ITYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[ITYPE:%.+]], i32 0, i32 0 +// CK1: [[IRES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[ISIZEADDR]], i64* [[ITYPEADDR]], i8** null) +// CK1: [[ISINITERR:%.+]] = icmp ne i32 [[IRES]], 0 +// CK1: br i1 [[ISINITERR]], label %[[INITERR:[^,]+]], label %[[LHEAD:[^,]+]] +// CK1: [[INITERR]] +// CK1: store i32 [[IRES]], i32* %retval +// CK1: br label %[[DONE:[^,]+]] + +// CK1: [[LHEAD]] +// CK1: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK1: br i1 [[ISEMPTY]], label %[[DONE]], label %[[LBODY:[^,]+]] +// CK1: [[LBODY]] +// CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK1: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] +// CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0 +// CK1-DAG: [[AEND:%.+]] = getelementptr i32, i32* [[ABEGIN]], i32 1 +// CK1-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8* +// CK1-DAG: [[AENDV:%.+]] = bitcast i32* [[AEND]] to i8* +// CK1-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64 +// CK1-DAG: [[AENDI:%.+]] = ptrtoint i8* [[AENDV]] to i64 +// CK1-DAG: [[CSIZE:%.+]] = sub i64 [[AENDI]], [[ABEGINI]] +// CK1-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK1-32-DAG: [[CUSIZE32:%.+]] = trunc i64 [[CUSIZE]] to i32 +// CK1-64-DAG: store i[[sz]] [[CUSIZE]], i[[sz]]* [[SIZEADDR0:%[^,]+]] +// CK1-32-DAG: store i[[sz]] [[CUSIZE32]], i[[sz]]* [[SIZEADDR0:%[^,]+]] +// CK1-DAG: [[SIZEADDR0]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR:%[^,]+]], i32 0, i32 0 +// CK1-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR0BC:%[^,]+]] +// CK1-DAG: [[BPTRADDR0BC]] = bitcast i8** [[BPTRADDR0:%.+]] to %class.C** +// CK1-DAG: [[BPTRADDR0]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR:%[^,]+]], i32 0, i32 0 +// CK1-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR0BC:%[^,]+]] +// CK1-DAG: [[PTRADDR0BC]] = bitcast i8** [[PTRADDR0:%.+]] to i32** +// CK1-DAG: [[PTRADDR0]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR:%[^,]+]], i32 0, i32 0 +// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK1-DAG: [[ALLOC]] +// CK1-DAG: br label %[[TYEND:[^,]+]] +// CK1-DAG: [[ALLOCELSE]] +// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK1-DAG: [[TO]] +// CK1-DAG: br label %[[TYEND]] +// CK1-DAG: [[TOELSE]] +// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK1-DAG: [[FROM]] +// CK1-DAG: br label %[[TYEND]] +// CK1-DAG: [[TYEND]] +// CK1-DAG: [[PHITYPE0:%.+]] = phi i64 [ 32, %[[ALLOC]] ], [ 32, %[[TO]] ], [ 32, %[[FROM]] ], [ 32, %[[TOELSE]] ] +// CK1-DAG: store i64 [[PHITYPE0]], i64* [[TYPEADDR0:%[^,]+]] +// CK1-DAG: [[TYPEADDR0]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR:%[^,]+]], i32 0, i32 0 +// CK1-DAG: [[BPTRADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR]], i32 0, i32 1 +// CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast i8** [[BPTRADDR1]] to %class.C** +// CK1-DAG: store %class.C* [[OBJ]], %class.C** [[BPTRADDR1BC]] +// CK1-DAG: [[PTRADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR]], i32 0, i32 1 +// CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i8** [[PTRADDR1]] to i32** +// CK1-DAG: store i32* [[ABEGIN]], i32** [[PTRADDR1BC]] +// CK1-DAG: [[SIZEADDR1:%.+]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 1 +// CK1-DAG: store i[[sz]] 4, i[[sz]]* [[SIZEADDR1]] +// CK1-DAG: [[TYPEADDR1:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR]], i32 0, i32 1 +// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK1-DAG: [[ALLOC]] +// CK1-DAG: br label %[[TYEND:[^,]+]] +// CK1-DAG: [[ALLOCELSE]] +// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK1-DAG: [[TO]] +// CK1-DAG: br label %[[TYEND]] +// CK1-DAG: [[TOELSE]] +// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK1-DAG: [[FROM]] +// CK1-DAG: br label %[[TYEND]] +// CK1-DAG: [[TYEND]] +// 0x1,000,000,000; 0x1,000,000,001; 0x1,000,000,002; 0x1,000,000,003 +// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ 281474976710656, %[[ALLOC]] ], [ 281474976710657, %[[TO]] ], [ 281474976710658, %[[FROM]] ], [ 281474976710659, %[[TOELSE]] ] +// CK1-DAG: store i64 [[TYPE1]], i64* [[TYPEADDR1]] +// CK1-DAG: [[ARGBPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[LBPTRADDR]], i32 0, i32 0 +// CK1-DAG: [[ARGPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRADDR]], i32 0, i32 0 +// CK1-DAG: [[ARGSIZE:%.+]] = getelementptr inbounds [2 x i[[sz]]], [2 x i[[sz]]]* [[SIZEADDR]], i32 0, i32 0 +// CK1-DAG: [[ARGTYPE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[TYPEADDR]], i32 0, i32 0 +// CK1: [[RES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 2, i8** [[ARGBPTR]], i8** [[ARGPTR]], i[[sz]]* [[ARGSIZE]], i64* [[ARGTYPE]], i8** null) +// CK1: [[ISERR:%.+]] = icmp ne i32 [[RES]], 0 +// CK1: br i1 [[ISERR]], label %[[LERR:[^,]+]], label %[[LCORRECT]] +// CK1: [[LERR]] +// CK1: store i32 [[RES]], i32* %retval +// CK1: br label %[[DONE]] +// CK1: [[LCORRECT]] +// CK1: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 +// CK1: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] + +// CK1: [[LEXIT]] +// CK1: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] +// CK1: [[EVALDEL]] +// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK1: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK1: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] +// CK1: [[DEL]] +// CK1: [[ARRSIZE:%.+]] = mul i[[sz]] [[SIZE]], 4 +// CK1-DAG: store i[[sz]] [[ARRSIZE]], i[[sz]]* [[DSIZEADDR:%[^,]+]] +// CK1-DAG: [[DSIZEADDR]] = getelementptr inbounds [1 x i[[sz]]], [1 x i[[sz]]]* [[DSIZE:%.+]], i32 0, i32 0 +// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK1-DAG: store i64 [[DTYPE]], i64* [[DTYPEADDR:%[^,]+]] +// CK1-DAG: [[DTYPEADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[DTYPE:%.+]], i32 0, i32 0 +// CK1: [[DRES:%.+]] = call i32 @__tgt_target_data_mapper_nowait(i64 [[DID]], i32 1, i8** [[BPTRADDR]], i8** [[VPTRADDR]], i[[sz]]* [[DSIZEADDR]], i64* [[DTYPEADDR]], i8** null) +// CK1: [[ISDELERR:%.+]] = icmp ne i32 [[DRES]], 0 +// CK1: br i1 [[ISDELERR]], label %[[DELERR:[^,]+]], label %[[DONE]] +// CK1: [[DELERR]] +// CK1: store i32 [[DRES]], i32* %retval +// CK1: br label %[[DONE]] +// CK1: [[DONE]] +// CK1: [[RET:%.+]] = load i32, i32* %retval +// CK1: ret i32 [[RET]] + +#endif #endif