From d985828d5afead32e20cb98d117c337f05cc32b9 Mon Sep 17 00:00:00 2001 From: Michael Kruse Date: Sun, 4 Aug 2019 05:16:52 +0000 Subject: [PATCH] Revert "[OpenMP 5.0] Codegen support for user-defined mappers." This reverts commit r367773. The test case OpenMP/declare_mapper_codegen.cpp is failing. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@367774 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/AST/GlobalDecl.h | 1 - lib/AST/ASTContext.cpp | 2 +- lib/CodeGen/CGDecl.cpp | 7 +- lib/CodeGen/CGOpenMPRuntime.cpp | 499 ++----------------------- lib/CodeGen/CGOpenMPRuntime.h | 20 - lib/CodeGen/ModuleBuilder.cpp | 3 - test/OpenMP/declare_mapper_codegen.cpp | 442 +++------------------- 7 files changed, 87 insertions(+), 887 deletions(-) diff --git a/include/clang/AST/GlobalDecl.h b/include/clang/AST/GlobalDecl.h index 145e961a23..86fd0f6aa9 100644 --- a/include/clang/AST/GlobalDecl.h +++ b/include/clang/AST/GlobalDecl.h @@ -59,7 +59,6 @@ public: 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) {} GlobalDecl(const VarDecl *D, DynamicInitKind StubKind) diff --git a/lib/AST/ASTContext.cpp b/lib/AST/ASTContext.cpp index be7830640e..468c7f4765 100644 --- a/lib/AST/ASTContext.cpp +++ b/lib/AST/ASTContext.cpp @@ -9860,7 +9860,7 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { return !D->getDeclContext()->isDependentContext(); else if (isa(D)) return !D->getDeclContext()->isDependentContext(); - else if (isa(D) || isa(D)) + else if (isa(D)) return !D->getDeclContext()->isDependentContext(); else if (isa(D)) return true; diff --git a/lib/CodeGen/CGDecl.cpp b/lib/CodeGen/CGDecl.cpp index a9e01619aa..c6c2ff57e1 100644 --- a/lib/CodeGen/CGDecl.cpp +++ b/lib/CodeGen/CGDecl.cpp @@ -2530,11 +2530,10 @@ void CodeGenModule::EmitOMPDeclareReduction(const OMPDeclareReductionDecl *D, } void CodeGenModule::EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D, - CodeGenFunction *CGF) { - if (!LangOpts.OpenMP || LangOpts.OpenMPSimd || - (!LangOpts.EmitAllDecls && !D->isUsed())) + CodeGenFunction *CGF) { + if (!LangOpts.OpenMP || (!LangOpts.EmitAllDecls && !D->isUsed())) return; - getOpenMPRuntime().emitUserDefinedMapper(D, CGF); + // FIXME: need to implement mapper code generation } void CodeGenModule::EmitOMPRequiresDecl(const OMPRequiresDecl *D) { diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp index 36b1a76cfc..2674e3a229 100644 --- a/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/lib/CodeGen/CGOpenMPRuntime.cpp @@ -752,11 +752,6 @@ enum OpenMPRTLFunction { // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t // *arg_types); OMPRTL__tgt_target_data_update_nowait, - // Call to int64_t __tgt_mapper_num_components(void *rt_mapper_handle); - OMPRTL__tgt_mapper_num_components, - // Call to void __tgt_push_mapper_component(void *rt_mapper_handle, void - // *base, void *begin, int64_t size, int64_t type); - OMPRTL__tgt_push_mapper_component, }; /// A basic class for pre|post-action for advanced codegen sequence for OpenMP @@ -1691,12 +1686,6 @@ void CGOpenMPRuntime::functionFinished(CodeGenFunction &CGF) { UDRMap.erase(D); FunctionUDRMap.erase(CGF.CurFn); } - auto I = FunctionUDMMap.find(CGF.CurFn); - if (I != FunctionUDMMap.end()) { - for(auto *D : I->second) - UDMMap.erase(D); - FunctionUDMMap.erase(I); - } } llvm::Type *CGOpenMPRuntime::getIdentTyPointerTy() { @@ -2470,24 +2459,6 @@ llvm::FunctionCallee CGOpenMPRuntime::createRuntimeFunction(unsigned Function) { RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_nowait"); break; } - case OMPRTL__tgt_mapper_num_components: { - // Build int64_t __tgt_mapper_num_components(void *rt_mapper_handle); - llvm::Type *TypeParams[] = {CGM.VoidPtrTy}; - auto *FnTy = - llvm::FunctionType::get(CGM.Int64Ty, TypeParams, /*isVarArg*/ false); - RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_mapper_num_components"); - break; - } - case OMPRTL__tgt_push_mapper_component: { - // Build void __tgt_push_mapper_component(void *rt_mapper_handle, void - // *base, void *begin, int64_t size, int64_t type); - llvm::Type *TypeParams[] = {CGM.VoidPtrTy, CGM.VoidPtrTy, CGM.VoidPtrTy, - CGM.Int64Ty, CGM.Int64Ty}; - auto *FnTy = - llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false); - RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_push_mapper_component"); - break; - } } assert(RTLFn && "Unable to find OpenMP runtime function"); return RTLFn; @@ -7122,15 +7093,6 @@ public: LLVM_MARK_AS_BITMASK_ENUM(/* LargestFlag = */ OMP_MAP_MEMBER_OF), }; - /// Get the offset of the OMP_MAP_MEMBER_OF field. - static unsigned getFlagMemberOffset() { - unsigned Offset = 0; - for (uint64_t Remain = OMP_MAP_MEMBER_OF; !(Remain & 1); - Remain = Remain >> 1) - Offset++; - return Offset; - } - /// Class that associates information with a base pointer to be passed to the /// runtime library. class BasePointerInfo { @@ -7194,11 +7156,8 @@ private: : IE(IE), VD(VD) {} }; - /// The target directive from where the mappable clauses were extracted. It - /// is either a executable directive or a user-defined mapper directive. - llvm::PointerUnion - CurDir; + /// Directive from where the map clauses were extracted. + const OMPExecutableDirective &CurDir; /// Function the directive is being generated for. CodeGenFunction &CGF; @@ -7801,9 +7760,9 @@ private: } static OpenMPOffloadMappingFlags getMemberOfFlag(unsigned Position) { - // Rotate by getFlagMemberOffset() bits. + // Member of is given by the 16 MSB of the flag, so rotate by 48 bits. return static_cast(((uint64_t)Position + 1) - << getFlagMemberOffset()); + << 48); } static void setCorrectMemberOfFlag(OpenMPOffloadMappingFlags &Flags, @@ -7883,7 +7842,7 @@ private: 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()) @@ -7895,10 +7854,6 @@ public: DevPointersMap[L.first].push_back(L.second); } - /// Constructor for the declare mapper directive. - MappableExprsHandler(const OMPDeclareMapperDecl &Dir, CodeGenFunction &CGF) - : CurDir(&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. @@ -7960,20 +7915,18 @@ public: IsImplicit); }; - assert(CurDir.is() && - "Expect a executable directive"); - const auto *CurExecDir = CurDir.get(); - for (const auto *C : CurExecDir->getClausesOfKind()) + // FIXME: MSVC 2013 seems to require this-> to find member CurDir. + 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 : CurExecDir->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 : CurExecDir->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()); @@ -7988,8 +7941,9 @@ public: llvm::MapVector> DeferredInfo; + // FIXME: MSVC 2013 seems to require this-> to find member CurDir. for (const auto *C : - CurExecDir->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(); @@ -8018,6 +7972,7 @@ public: // We didn't find any match in our map information - generate a zero // size array section - if the pointer is a struct member we defer this // action until the whole struct has been processed. + // FIXME: MSVC 2013 seems to require this-> to find member CGF. if (isa(IE)) { // Insert the pointer into Info to be processed by // generateInfoForComponentList. Because it is a member pointer @@ -8030,11 +7985,11 @@ public: /*ReturnDevicePointer=*/false, C->isImplicit()); DeferredInfo[nullptr].emplace_back(IE, VD); } else { - llvm::Value *Ptr = - CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc()); + llvm::Value *Ptr = this->CGF.EmitLoadOfScalar( + this->CGF.EmitLValue(IE), IE->getExprLoc()); BasePointers.emplace_back(Ptr, VD); Pointers.push_back(Ptr); - Sizes.push_back(llvm::Constant::getNullValue(CGF.Int64Ty)); + Sizes.push_back(llvm::Constant::getNullValue(this->CGF.Int64Ty)); Types.push_back(OMP_MAP_RETURN_PARAM | OMP_MAP_TARGET_PARAM); } } @@ -8058,10 +8013,11 @@ public: // Remember the current base pointer index. unsigned CurrentBasePointersIdx = CurBasePointers.size(); - generateInfoForComponentList(L.MapType, L.MapModifiers, L.Components, - CurBasePointers, CurPointers, CurSizes, - CurTypes, PartialStruct, - IsFirstComponentList, L.IsImplicit); + // 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); // If this entry relates with a device pointer, set the relevant // declaration and add the 'return pointer' flag. @@ -8113,78 +8069,6 @@ public: } } - /// Generate all the base pointers, section pointers, sizes and map types for - /// the extracted map clauses of user-defined mapper. - void generateAllInfoForMapper(MapBaseValuesArrayTy &BasePointers, - MapValuesArrayTy &Pointers, - MapValuesArrayTy &Sizes, - MapFlagsArrayTy &Types) const { - assert(CurDir.is() && - "Expect a declare mapper directive"); - const auto *CurMapperDir = CurDir.get(); - // 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); - }; - - for (const auto *C : 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."); - 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, @@ -8308,10 +8192,8 @@ public: std::tuple, bool>; SmallVector DeclComponentLists; - assert(CurDir.is() && - "Expect a executable directive"); - const auto *CurExecDir = CurDir.get(); - for (const auto *C : CurExecDir->getClausesOfKind()) { + // FIXME: MSVC 2013 seems to require this-> to find member CurDir. + 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??"); @@ -8459,12 +8341,9 @@ public: MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types) const { - assert(CurDir.is() && - "Expect a executable directive"); - const auto *CurExecDir = CurDir.get(); // Map other list items in the map clause which are not captured variables // but "declare target link" global variables. - for (const auto *C : CurExecDir->getClausesOfKind()) { + for (const auto *C : this->CurDir.getClausesOfKind()) { for (const auto &L : C->component_lists()) { if (!L.first) continue; @@ -8691,7 +8570,6 @@ emitOffloadingArrays(CodeGenFunction &CGF, } } } - /// Emit the arguments to be passed to the runtime library based on the /// arrays of pointers, sizes and map types. static void emitOffloadingArraysArgument( @@ -8822,337 +8700,6 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) { return nullptr; } -/// Emit the user-defined mapper function. The code generation follows the -/// pattern in the example below. -/// \code -/// void .omp_mapper...(void *rt_mapper_handle, -/// void *base, void *begin, -/// int64_t size, int64_t type) { -/// // Allocate space for an array section first. -/// if (size > 1 && !maptype.IsDelete) -/// __tgt_push_mapper_component(rt_mapper_handle, base, begin, -/// size*sizeof(Ty), clearToFrom(type)); -/// // Map members. -/// for (unsigned i = 0; i < size; i++) { -/// // For each component specified by this mapper: -/// for (auto c : all_components) { -/// if (c.hasMapper()) -/// (*c.Mapper())(rt_mapper_handle, c.arg_base, c.arg_begin, c.arg_size, -/// c.arg_type); -/// else -/// __tgt_push_mapper_component(rt_mapper_handle, c.arg_base, -/// c.arg_begin, c.arg_size, c.arg_type); -/// } -/// } -/// // Delete the array section. -/// if (size > 1 && maptype.IsDelete) -/// __tgt_push_mapper_component(rt_mapper_handle, base, begin, -/// size*sizeof(Ty), clearToFrom(type)); -/// } -/// \endcode -void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D, - CodeGenFunction *CGF) { - if (UDMMap.count(D) > 0) - return; - ASTContext &C = CGM.getContext(); - QualType Ty = D->getType(); - QualType PtrTy = C.getPointerType(Ty).withRestrict(); - 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 HandleArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, - C.VoidPtrTy, ImplicitParamDecl::Other); - ImplicitParamDecl BaseArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.VoidPtrTy, - ImplicitParamDecl::Other); - ImplicitParamDecl BeginArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, - C.VoidPtrTy, ImplicitParamDecl::Other); - ImplicitParamDecl SizeArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, Int64Ty, - ImplicitParamDecl::Other); - ImplicitParamDecl TypeArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, Int64Ty, - ImplicitParamDecl::Other); - FunctionArgList Args; - Args.push_back(&HandleArg); - Args.push_back(&BaseArg); - Args.push_back(&BeginArg); - Args.push_back(&SizeArg); - Args.push_back(&TypeArg); - const CGFunctionInfo &FnInfo = - CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args); - llvm::FunctionType *FnTy = CGM.getTypes().GetFunctionType(FnInfo); - SmallString<64> TyStr; - llvm::raw_svector_ostream Out(TyStr); - CGM.getCXXABI().getMangleContext().mangleTypeName(Ty, Out); - std::string Name = getName({"omp_mapper", TyStr, D->getName()}); - 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.VoidTy, Fn, FnInfo, Args, Loc, Loc); - // Compute the starting and end addreses of array elements. - llvm::Value *Size = MapperCGF.EmitLoadOfScalar( - MapperCGF.GetAddrOfLocalVar(&SizeArg), /*Volatile=*/false, - C.getPointerType(Int64Ty), Loc); - llvm::Value *PtrBegin = MapperCGF.Builder.CreateBitCast( - MapperCGF.GetAddrOfLocalVar(&BeginArg).getPointer(), - CGM.getTypes().ConvertTypeForMem(C.getPointerType(PtrTy))); - llvm::Value *PtrEnd = MapperCGF.Builder.CreateGEP(PtrBegin, Size); - llvm::Value *MapType = MapperCGF.EmitLoadOfScalar( - MapperCGF.GetAddrOfLocalVar(&TypeArg), /*Volatile=*/false, - C.getPointerType(Int64Ty), Loc); - // Prepare common arguments for array initiation and deletion. - llvm::Value *Handle = MapperCGF.EmitLoadOfScalar( - MapperCGF.GetAddrOfLocalVar(&HandleArg), - /*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc); - llvm::Value *BaseIn = MapperCGF.EmitLoadOfScalar( - MapperCGF.GetAddrOfLocalVar(&BaseArg), - /*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc); - llvm::Value *BeginIn = MapperCGF.EmitLoadOfScalar( - MapperCGF.GetAddrOfLocalVar(&BeginArg), - /*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc); - - // Emit array initiation if this is an array section and \p MapType indicates - // that memory allocation is required. - llvm::BasicBlock *HeadBB = MapperCGF.createBasicBlock("omp.arraymap.head"); - emitUDMapperArrayInitOrDel(MapperCGF, Handle, BaseIn, BeginIn, Size, MapType, - ElementSize, HeadBB, /*IsInit=*/true); - - // 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 *DoneBB = MapperCGF.createBasicBlock("omp.done"); - // 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(&BeginArg) - .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 mapped variables. - MappableExprsHandler::MapBaseValuesArrayTy BasePointers; - MappableExprsHandler::MapValuesArrayTy Pointers; - MappableExprsHandler::MapValuesArrayTy Sizes; - MappableExprsHandler::MapFlagsArrayTy MapTypes; - MappableExprsHandler MEHandler(*D, MapperCGF); - MEHandler.generateAllInfoForMapper(BasePointers, Pointers, Sizes, MapTypes); - - // Call the runtime API __tgt_mapper_num_components to get the number of - // pre-existing components. - llvm::Value *OffloadingArgs[] = {Handle}; - llvm::Value *PreviousSize = MapperCGF.EmitRuntimeCall( - createRuntimeFunction(OMPRTL__tgt_mapper_num_components), OffloadingArgs); - llvm::Value *ShiftedPreviousSize = MapperCGF.Builder.CreateShl( - PreviousSize, - MapperCGF.Builder.getInt64(MappableExprsHandler::getFlagMemberOffset())); - - // Fill up the runtime mapper handle for all components. - for (unsigned I = 0; I < BasePointers.size(); ++I) { - llvm::Value *CurBaseArg = MapperCGF.Builder.CreateBitCast( - *BasePointers[I], CGM.getTypes().ConvertTypeForMem(C.VoidPtrTy)); - llvm::Value *CurBeginArg = MapperCGF.Builder.CreateBitCast( - Pointers[I], CGM.getTypes().ConvertTypeForMem(C.VoidPtrTy)); - llvm::Value *CurSizeArg = Sizes[I]; - - // Extract the MEMBER_OF field from the map type. - llvm::BasicBlock *MemberBB = MapperCGF.createBasicBlock("omp.member"); - MapperCGF.EmitBlock(MemberBB); - llvm::Value *OriMapType = MapperCGF.Builder.getInt64(MapTypes[I]); - llvm::Value *Member = MapperCGF.Builder.CreateAnd( - OriMapType, - MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_MEMBER_OF)); - llvm::BasicBlock *MemberCombineBB = - MapperCGF.createBasicBlock("omp.member.combine"); - llvm::BasicBlock *TypeBB = MapperCGF.createBasicBlock("omp.type"); - llvm::Value *IsMember = MapperCGF.Builder.CreateIsNull(Member); - MapperCGF.Builder.CreateCondBr(IsMember, TypeBB, MemberCombineBB); - // Add the number of pre-existing components to the MEMBER_OF field if it - // is valid. - MapperCGF.EmitBlock(MemberCombineBB); - llvm::Value *CombinedMember = - MapperCGF.Builder.CreateNUWAdd(OriMapType, ShiftedPreviousSize); - // Do nothing if it is not a member of previous components. - MapperCGF.EmitBlock(TypeBB); - llvm::PHINode *MemberMapType = - MapperCGF.Builder.CreatePHI(CGM.Int64Ty, 4, "omp.membermaptype"); - MemberMapType->addIncoming(OriMapType, MemberBB); - MemberMapType->addIncoming(CombinedMember, MemberCombineBB); - - // Combine the map type inherited from user-defined mapper with that - // specified in the program. According to the OMP_MAP_TO and OMP_MAP_FROM - // bits of the \a MapType, which is the input argument of the mapper - // function, the following code will set the OMP_MAP_TO and OMP_MAP_FROM - // bits of MemberMapType. - // [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 *LeftToFrom = MapperCGF.Builder.CreateAnd( - MapType, - MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_TO | - MappableExprsHandler::OMP_MAP_FROM)); - llvm::BasicBlock *AllocBB = MapperCGF.createBasicBlock("omp.type.alloc"); - llvm::BasicBlock *AllocElseBB = - MapperCGF.createBasicBlock("omp.type.alloc.else"); - llvm::BasicBlock *ToBB = MapperCGF.createBasicBlock("omp.type.to"); - llvm::BasicBlock *ToElseBB = MapperCGF.createBasicBlock("omp.type.to.else"); - llvm::BasicBlock *FromBB = MapperCGF.createBasicBlock("omp.type.from"); - llvm::BasicBlock *EndBB = MapperCGF.createBasicBlock("omp.type.end"); - llvm::Value *IsAlloc = MapperCGF.Builder.CreateIsNull(LeftToFrom); - MapperCGF.Builder.CreateCondBr(IsAlloc, AllocBB, AllocElseBB); - // In case of alloc, clear OMP_MAP_TO and OMP_MAP_FROM. - MapperCGF.EmitBlock(AllocBB); - llvm::Value *AllocMapType = MapperCGF.Builder.CreateAnd( - MemberMapType, - MapperCGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO | - MappableExprsHandler::OMP_MAP_FROM))); - MapperCGF.Builder.CreateBr(EndBB); - MapperCGF.EmitBlock(AllocElseBB); - llvm::Value *IsTo = MapperCGF.Builder.CreateICmpEQ( - LeftToFrom, - MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_TO)); - MapperCGF.Builder.CreateCondBr(IsTo, ToBB, ToElseBB); - // In case of to, clear OMP_MAP_FROM. - MapperCGF.EmitBlock(ToBB); - llvm::Value *ToMapType = MapperCGF.Builder.CreateAnd( - MemberMapType, - MapperCGF.Builder.getInt64(~MappableExprsHandler::OMP_MAP_FROM)); - MapperCGF.Builder.CreateBr(EndBB); - MapperCGF.EmitBlock(ToElseBB); - llvm::Value *IsFrom = MapperCGF.Builder.CreateICmpEQ( - LeftToFrom, - MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_FROM)); - MapperCGF.Builder.CreateCondBr(IsFrom, FromBB, EndBB); - // In case of from, clear OMP_MAP_TO. - MapperCGF.EmitBlock(FromBB); - llvm::Value *FromMapType = MapperCGF.Builder.CreateAnd( - MemberMapType, - MapperCGF.Builder.getInt64(~MappableExprsHandler::OMP_MAP_TO)); - // In case of tofrom, do nothing. - MapperCGF.EmitBlock(EndBB); - llvm::PHINode *CurMapType = - MapperCGF.Builder.CreatePHI(CGM.Int64Ty, 4, "omp.maptype"); - CurMapType->addIncoming(AllocMapType, AllocBB); - CurMapType->addIncoming(ToMapType, ToBB); - CurMapType->addIncoming(FromMapType, FromBB); - CurMapType->addIncoming(MemberMapType, ToElseBB); - - // TODO: call the corresponding mapper function if a user-defined mapper is - // associated with this map clause. - // Call the runtime API __tgt_push_mapper_component to fill up the runtime - // data structure. - llvm::Value *OffloadingArgs[] = {Handle, CurBaseArg, CurBeginArg, - CurSizeArg, CurMapType}; - MapperCGF.EmitRuntimeCall( - createRuntimeFunction(OMPRTL__tgt_push_mapper_component), - OffloadingArgs); - } - - // Update the pointer to point to the next element that needs to be mapped, - // and check whether we have mapped all elements. - llvm::Value *PtrNext = MapperCGF.Builder.CreateConstGEP1_32( - PtrPHI, /*Idx0=*/1, "omp.arraymap.next"); - PtrPHI->addIncoming(PtrNext, BodyBB); - llvm::Value *IsDone = - MapperCGF.Builder.CreateICmpEQ(PtrNext, PtrEnd, "omp.arraymap.isdone"); - llvm::BasicBlock *ExitBB = MapperCGF.createBasicBlock("omp.arraymap.exit"); - MapperCGF.Builder.CreateCondBr(IsDone, ExitBB, BodyBB); - - MapperCGF.EmitBlock(ExitBB); - // Emit array deletion if this is an array section and \p MapType indicates - // that deletion is required. - emitUDMapperArrayInitOrDel(MapperCGF, Handle, BaseIn, BeginIn, Size, MapType, - ElementSize, DoneBB, /*IsInit=*/false); - - // Emit the function exit block. - MapperCGF.EmitBlock(DoneBB, /*IsFinished=*/true); - MapperCGF.FinishFunction(); - UDMMap.try_emplace(D, Fn); - if (CGF) { - auto &Decls = FunctionUDMMap.FindAndConstruct(CGF->CurFn); - Decls.second.push_back(D); - } -} - -/// Emit the array initialization or deletion portion for user-defined mapper -/// code generation. First, it evaluates whether an array section is mapped and -/// whether the \a MapType instructs to delete this section. If \a IsInit is -/// true, and \a MapType indicates to not delete this array, array -/// initialization code is generated. If \a IsInit is false, and \a MapType -/// indicates to not this array, array deletion code is generated. -void CGOpenMPRuntime::emitUDMapperArrayInitOrDel( - CodeGenFunction &MapperCGF, llvm::Value *Handle, llvm::Value *Base, - llvm::Value *Begin, llvm::Value *Size, llvm::Value *MapType, - CharUnits ElementSize, llvm::BasicBlock *ExitBB, bool IsInit) { - StringRef Prefix = IsInit ? ".init" : ".del"; - - // Evaluate if this is an array section. - llvm::BasicBlock *IsDeleteBB = - MapperCGF.createBasicBlock("omp.array" + Prefix + ".evaldelete"); - llvm::BasicBlock *BodyBB = MapperCGF.createBasicBlock("omp.array" + Prefix); - llvm::Value *IsArray = MapperCGF.Builder.CreateICmpSGE( - Size, MapperCGF.Builder.getInt64(1), "omp.arrayinit.isarray"); - MapperCGF.Builder.CreateCondBr(IsArray, IsDeleteBB, ExitBB); - - // Evaluate if we are going to delete this section. - MapperCGF.EmitBlock(IsDeleteBB); - llvm::Value *DeleteBit = MapperCGF.Builder.CreateAnd( - MapType, - MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_DELETE)); - llvm::Value *DeleteCond; - if (IsInit) { - DeleteCond = MapperCGF.Builder.CreateIsNull( - DeleteBit, "omp.array" + Prefix + ".delete"); - } else { - DeleteCond = MapperCGF.Builder.CreateIsNotNull( - DeleteBit, "omp.array" + Prefix + ".delete"); - } - MapperCGF.Builder.CreateCondBr(DeleteCond, BodyBB, ExitBB); - - MapperCGF.EmitBlock(BodyBB); - // Get the array size by multiplying element size and element number (i.e., \p - // Size). - llvm::Value *ArraySize = MapperCGF.Builder.CreateNUWMul( - Size, MapperCGF.Builder.getInt64(ElementSize.getQuantity())); - // Remove OMP_MAP_TO and OMP_MAP_FROM from the map type, so that it achieves - // memory allocation/deletion purpose only. - llvm::Value *MapTypeArg = MapperCGF.Builder.CreateAnd( - MapType, - MapperCGF.Builder.getInt64(~(MappableExprsHandler::OMP_MAP_TO | - MappableExprsHandler::OMP_MAP_FROM))); - // Call the runtime API __tgt_push_mapper_component to fill up the runtime - // data structure. - llvm::Value *OffloadingArgs[] = {Handle, Base, Begin, ArraySize, MapTypeArg}; - MapperCGF.EmitRuntimeCall( - createRuntimeFunction(OMPRTL__tgt_push_mapper_component), OffloadingArgs); -} - void CGOpenMPRuntime::emitTargetNumIterationsCall( CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *Device, const llvm::function_ref> FunctionUDRMapTy; FunctionUDRMapTy FunctionUDRMap; - /// Map from the user-defined mapper declaration to its corresponding - /// functions. - llvm::DenseMap UDMMap; - /// Map of functions and their local user-defined mappers. - using FunctionUDMMapTy = - llvm::DenseMap>; - FunctionUDMMapTy FunctionUDMMap; /// Type kmp_critical_name, originally defined as typedef kmp_int32 /// kmp_critical_name[8]; llvm::ArrayType *KmpCriticalNameTy; @@ -746,14 +738,6 @@ private: llvm::Value *Ctor, llvm::Value *CopyCtor, llvm::Value *Dtor, SourceLocation Loc); - /// Emit the array initialization or deletion portion for user-defined mapper - /// code generation. - void emitUDMapperArrayInitOrDel(CodeGenFunction &MapperCGF, - llvm::Value *Handle, llvm::Value *BasePtr, - llvm::Value *Ptr, llvm::Value *Size, - llvm::Value *MapType, CharUnits ElementSize, - llvm::BasicBlock *ExitBB, bool IsInit); - struct TaskResultTy { llvm::Value *NewTask = nullptr; llvm::Function *TaskEntry = nullptr; @@ -814,10 +798,6 @@ public: virtual std::pair getUserDefinedReduction(const OMPDeclareReductionDecl *D); - /// Emit the function for the user defined mapper construct. - void emitUserDefinedMapper(const OMPDeclareMapperDecl *D, - CodeGenFunction *CGF = nullptr); - /// 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*). diff --git a/lib/CodeGen/ModuleBuilder.cpp b/lib/CodeGen/ModuleBuilder.cpp index 414a6b8976..3b4e06045a 100644 --- a/lib/CodeGen/ModuleBuilder.cpp +++ b/lib/CodeGen/ModuleBuilder.cpp @@ -232,9 +232,6 @@ namespace { if (auto *DRD = dyn_cast(Member)) { if (Ctx->DeclMustBeEmitted(DRD)) Builder->EmitGlobal(DRD); - } else if (auto *DMD = dyn_cast(Member)) { - if (Ctx->DeclMustBeEmitted(DMD)) - Builder->EmitGlobal(DMD); } } } diff --git a/test/OpenMP/declare_mapper_codegen.cpp b/test/OpenMP/declare_mapper_codegen.cpp index dc9e26affc..6f1d6ec8fd 100644 --- a/test/OpenMP/declare_mapper_codegen.cpp +++ b/test/OpenMP/declare_mapper_codegen.cpp @@ -1,414 +1,92 @@ +///==========================================================================/// +// 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 i64] [i64 16] -// CK0-32: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8] -// CK0: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35] -// CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] -// CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8] -// CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33] -// CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] -// CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 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, s.b[0:2]) +#pragma omp declare mapper(id: C s) map(s.a) -// CK0-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C.id{{.*}}(i8*, i8*, i8*, i64, i64) -// CK0: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]] -// CK0: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]] -// CK0: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]] -// CK0: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]] -// CK0: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]] -// CK0-DAG: [[SIZE:%.+]] = load i64, i64* [[SIZEADDR]] -// CK0-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] -// CK0-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]] -// CK0-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** -// CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]] -// CK0-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]] -// CK0-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]] -// CK0: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// CK0: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] +// CHECK-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l54.region_id = weak constant i8 0 -// 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-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 -// CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 -// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 -// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]]) -// CK0: br label %[[LHEAD:[^,]+]] +// 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] -// 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:64|32]] 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-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* -// CK0-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* -// CK0-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) -// CK0-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 -// CK0-DAG: br label %[[MEMBER:[^,]+]] -// CK0-DAG: [[MEMBER]] -// CK0-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK0-DAG: [[MEMBERCOM]] -// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]] -// CK0-DAG: br label %[[LTYPE]] -// CK0-DAG: [[LTYPE]] -// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] -// 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: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 -// 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: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 -// 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: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 -// CK0-DAG: br label %[[TYEND]] -// CK0-DAG: [[TYEND]] -// CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]]) -// CK0-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* -// CK0-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* -// CK0-DAG: br label %[[MEMBER:[^,]+]] -// CK0-DAG: [[MEMBER]] -// CK0-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK0-DAG: [[MEMBERCOM]] -// 281474976710659 == 0x1,000,000,003 -// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] -// CK0-DAG: br label %[[LTYPE]] -// CK0-DAG: [[LTYPE]] -// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] -// 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: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 -// 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: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 -// 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: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 -// CK0-DAG: br label %[[TYEND]] -// CK0-DAG: [[TYEND]] -// CK0-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]]) -// CK0-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8* -// CK0-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8* -// CK0-DAG: br label %[[MEMBER:[^,]+]] -// CK0-DAG: [[MEMBER]] -// CK0-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK0-DAG: [[MEMBERCOM]] -// 281474976710675 == 0x1,000,000,013 -// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]] -// CK0-DAG: br label %[[LTYPE]] -// CK0-DAG: [[LTYPE]] -// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710675, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] -// 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: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 -// 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: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 -// 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: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 -// CK0-DAG: br label %[[TYEND]] -// CK0-DAG: [[TYEND]] -// CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]]) -// 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: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// 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-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 -// CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 -// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 -// CK0: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]]) -// CK0: br label %[[DONE]] -// CK0: [[DONE]] -// CK0: ret void - - -// CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}} +// CHECK-LABEL: foo{{.*}}( void foo(int a){ int i = a; C c; c.a = a; - // 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]]) + // 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]]) #pragma omp target map(mapper(id),tofrom: c) { ++c.a; } - // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[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]] + // 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]] #pragma omp target update to(mapper(id): c) - // CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[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]] + // 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]] #pragma omp target update from(mapper(id): c) } -// 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) - -// CK1-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id{{.*}}(i8*, i8*, i8*, i64, i64) -// CK1: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]] -// CK1: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]] -// CK1: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]] -// CK1: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]] -// CK1: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]] -// CK1-DAG: [[SIZE:%.+]] = load i64, i64* [[SIZEADDR]] -// CK1-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] -// CK1-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]] -// CK1-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** -// CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]] -// CK1-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]] -// CK1-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]] -// CK1: [[ISARRAY:%.+]] = icmp sge i64 [[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-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4 -// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 -// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]]) -// CK1: br label %[[LHEAD:[^,]+]] - -// 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-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* -// CK1-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* -// CK1-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) -// CK1-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 -// CK1-DAG: br label %[[MEMBER:[^,]+]] -// CK1-DAG: [[MEMBER]] -// CK1-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK1-DAG: [[MEMBERCOM]] -// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]] -// CK1-DAG: br label %[[LTYPE]] -// CK1-DAG: [[LTYPE]] -// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] -// 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: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 -// 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: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 -// 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: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 -// CK1-DAG: br label %[[TYEND]] -// CK1-DAG: [[TYEND]] -// CK1-DAG: [[TYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK1-64: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[TYPE0]]) -// CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* -// CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* -// CK1-DAG: br label %[[MEMBER:[^,]+]] -// CK1-DAG: [[MEMBER]] -// CK1-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK1-DAG: [[MEMBERCOM]] -// 281474976710659 == 0x1,000,000,003 -// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] -// CK1-DAG: br label %[[LTYPE]] -// CK1-DAG: [[LTYPE]] -// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] -// 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: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 -// 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: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 -// 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: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 -// CK1-DAG: br label %[[TYEND]] -// CK1-DAG: [[TYEND]] -// CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]]) -// 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: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 -// 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-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4 -// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 -// CK1: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]]) -// CK1: br label %[[DONE]] -// CK1: [[DONE]] -// CK1: ret void - -#endif +// 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: } #endif -- 2.40.0