Index: include/clang/AST/OpenMPClause.h =================================================================== --- include/clang/AST/OpenMPClause.h +++ include/clang/AST/OpenMPClause.h @@ -4286,6 +4286,11 @@ /// Total number of components in this clause. unsigned NumComponents; + /// Whether this clause is possible to have user-defined mappers associated. + /// It should be true for map, to, and from clauses, and false for + /// use_device_ptr and is_device_ptr. + bool hasMapper; + /// C++ nested name specifier for the associated user-defined mapper. NestedNameSpecifierLoc MapperQualifierLoc; @@ -4306,19 +4311,21 @@ /// NumUniqueDeclarations: number of unique base declarations in this clause; /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. + /// \param hasMapper Indicates whether this clause is possible to have + /// user-defined mappers associated. /// \param MapperQualifierLocPtr C++ nested name specifier for the associated /// user-defined mapper. /// \param MapperIdInfoPtr The identifier of associated user-defined mapper. OMPMappableExprListClause( OpenMPClauseKind K, const OMPVarListLocTy &Locs, - const OMPMappableExprListSizeTy &Sizes, + const OMPMappableExprListSizeTy &Sizes, bool hasMapper = false, NestedNameSpecifierLoc *MapperQualifierLocPtr = nullptr, DeclarationNameInfo *MapperIdInfoPtr = nullptr) : OMPVarListClause(K, Locs.StartLoc, Locs.LParenLoc, Locs.EndLoc, Sizes.NumVars), NumUniqueDeclarations(Sizes.NumUniqueDeclarations), NumComponentLists(Sizes.NumComponentLists), - NumComponents(Sizes.NumComponents) { + NumComponents(Sizes.NumComponents), hasMapper(hasMapper) { if (MapperQualifierLocPtr) MapperQualifierLoc = *MapperQualifierLocPtr; if (MapperIdInfoPtr) @@ -4517,6 +4524,8 @@ /// Get the user-defined mapper references that are in the trailing objects of /// the class. MutableArrayRef getUDMapperRefs() { + assert(hasMapper && + "Must be a clause that is possible to have user-defined mappers"); return llvm::makeMutableArrayRef( static_cast(this)->template getTrailingObjects() + OMPVarListClause::varlist_size(), @@ -4525,9 +4534,11 @@ /// Get the user-defined mappers references that are in the trailing objects /// of the class. - ArrayRef getUDMapperRefs() const { - return llvm::makeArrayRef( - static_cast(this)->template getTrailingObjects() + + ArrayRef getUDMapperRefs() const { + assert(hasMapper && + "Must be a clause that is possible to have user-defined mappers"); + return llvm::makeArrayRef( + static_cast(this)->template getTrailingObjects() + OMPVarListClause::varlist_size(), OMPVarListClause::varlist_size()); } @@ -4537,6 +4548,8 @@ void setUDMapperRefs(ArrayRef DMDs) { assert(DMDs.size() == OMPVarListClause::varlist_size() && "Unexpected number of user-defined mappers."); + assert(hasMapper && + "Must be a clause that is possible to have user-defined mappers"); std::copy(DMDs.begin(), DMDs.end(), getUDMapperRefs().begin()); } @@ -4573,6 +4586,12 @@ // The list number associated with the current declaration. ArrayRef::iterator NumListsCur; + // Whether this clause is possible to have user-defined mappers associated. + bool hasMapper; + + // The user-defined mapper associated with the current declaration. + ArrayRef::iterator MapperCur; + // Remaining lists for the current declaration. unsigned RemainingLists = 0; @@ -4593,10 +4612,12 @@ explicit const_component_lists_iterator( ArrayRef UniqueDecls, ArrayRef DeclsListNum, ArrayRef CumulativeListSizes, - MappableExprComponentListRef Components) + MappableExprComponentListRef Components, bool hasMapper, + ArrayRef Mappers) : const_component_lists_iterator::iterator_adaptor_base( Components.begin()), DeclCur(UniqueDecls.begin()), NumListsCur(DeclsListNum.begin()), + hasMapper(hasMapper), MapperCur(Mappers.begin()), ListSizeCur(CumulativeListSizes.begin()), ListSizeEnd(CumulativeListSizes.end()), End(Components.end()) { assert(UniqueDecls.size() == DeclsListNum.size() && @@ -4610,9 +4631,11 @@ explicit const_component_lists_iterator( const ValueDecl *Declaration, ArrayRef UniqueDecls, ArrayRef DeclsListNum, ArrayRef CumulativeListSizes, - MappableExprComponentListRef Components) + MappableExprComponentListRef Components, bool hasMapper, + ArrayRef Mappers) : const_component_lists_iterator(UniqueDecls, DeclsListNum, - CumulativeListSizes, Components) { + CumulativeListSizes, Components, + hasMapper, Mappers) { // Look for the desired declaration. While we are looking for it, we // update the state so that we know the component where a given list // starts. @@ -4627,6 +4650,9 @@ std::advance(ListSizeCur, *NumListsCur - 1); PrevListSize = *ListSizeCur; ++ListSizeCur; + + if (hasMapper) + ++MapperCur; } // If we didn't find any declaration, advance the iterator to after the @@ -4652,14 +4678,20 @@ // Return the array with the current list. The sizes are cumulative, so the // array size is the difference between the current size and previous one. - std::pair + std::tuple operator*() const { assert(ListSizeCur != ListSizeEnd && "Invalid iterator!"); - return std::make_pair( + const ValueDecl *Mapper = nullptr; + if (hasMapper && *MapperCur) + Mapper = cast(cast(*MapperCur)->getDecl()); + return std::make_tuple( *DeclCur, - MappableExprComponentListRef(&*this->I, *ListSizeCur - PrevListSize)); + MappableExprComponentListRef(&*this->I, *ListSizeCur - PrevListSize), + Mapper); } - std::pair + std::tuple operator->() const { return **this; } @@ -4682,6 +4714,8 @@ if (!(--RemainingLists)) { ++DeclCur; ++NumListsCur; + if (hasMapper) + ++MapperCur; RemainingLists = *NumListsCur; assert(RemainingLists && "No lists in the following declaration??"); } @@ -4699,13 +4733,15 @@ const_component_lists_iterator component_lists_begin() const { return const_component_lists_iterator( getUniqueDeclsRef(), getDeclNumListsRef(), getComponentListSizesRef(), - getComponentsRef()); + getComponentsRef(), hasMapper, + hasMapper ? getUDMapperRefs() : ArrayRef()); } const_component_lists_iterator component_lists_end() const { return const_component_lists_iterator( ArrayRef(), ArrayRef(), ArrayRef(), MappableExprComponentListRef(getComponentsRef().end(), - getComponentsRef().end())); + getComponentsRef().end()), + hasMapper, ArrayRef()); } const_component_lists_range component_lists() const { return {component_lists_begin(), component_lists_end()}; @@ -4717,7 +4753,8 @@ decl_component_lists_begin(const ValueDecl *VD) const { return const_component_lists_iterator( VD, getUniqueDeclsRef(), getDeclNumListsRef(), - getComponentListSizesRef(), getComponentsRef()); + getComponentListSizesRef(), getComponentsRef(), hasMapper, + hasMapper ? getUDMapperRefs() : ArrayRef()); } const_component_lists_iterator decl_component_lists_end() const { return component_lists_end(); @@ -4869,8 +4906,8 @@ OpenMPMapClauseKind MapType, bool MapTypeIsImplicit, SourceLocation MapLoc, const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(OMPC_map, Locs, Sizes, &MapperQualifierLoc, - &MapperIdInfo), + : OMPMappableExprListClause(OMPC_map, Locs, Sizes, /*hasMapper=*/true, + &MapperQualifierLoc, &MapperIdInfo), MapType(MapType), MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) { assert(llvm::array_lengthof(MapTypeModifiers) == MapModifiers.size() && "Unexpected number of map type modifiers."); @@ -4890,7 +4927,8 @@ /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. explicit OMPMapClause(const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(OMPC_map, OMPVarListLocTy(), Sizes) {} + : OMPMappableExprListClause(OMPC_map, OMPVarListLocTy(), Sizes, + /*hasMapper=*/true) {} /// Set map-type-modifier for the clause. /// @@ -5744,8 +5782,8 @@ DeclarationNameInfo MapperIdInfo, const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(OMPC_to, Locs, Sizes, &MapperQualifierLoc, - &MapperIdInfo) {} + : OMPMappableExprListClause(OMPC_to, Locs, Sizes, /*hasMapper=*/true, + &MapperQualifierLoc, &MapperIdInfo) {} /// Build an empty clause. /// @@ -5755,7 +5793,8 @@ /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. explicit OMPToClause(const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(OMPC_to, OMPVarListLocTy(), Sizes) {} + : OMPMappableExprListClause(OMPC_to, OMPVarListLocTy(), Sizes, + /*hasMapper=*/true) {} /// Define the sizes of each trailing object array except the last one. This /// is required for TrailingObjects to work properly. @@ -5862,8 +5901,8 @@ DeclarationNameInfo MapperIdInfo, const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(OMPC_from, Locs, Sizes, &MapperQualifierLoc, - &MapperIdInfo) {} + : OMPMappableExprListClause(OMPC_from, Locs, Sizes, /*hasMapper=*/true, + &MapperQualifierLoc, &MapperIdInfo) {} /// Build an empty clause. /// @@ -5873,7 +5912,8 @@ /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. explicit OMPFromClause(const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(OMPC_from, OMPVarListLocTy(), Sizes) {} + : OMPMappableExprListClause(OMPC_from, OMPVarListLocTy(), Sizes, + /*hasMapper=*/true) {} /// Define the sizes of each trailing object array except the last one. This /// is required for TrailingObjects to work properly. Index: lib/CodeGen/CGOpenMPRuntime.h =================================================================== --- lib/CodeGen/CGOpenMPRuntime.h +++ lib/CodeGen/CGOpenMPRuntime.h @@ -817,6 +817,9 @@ /// Emit the function for the user defined mapper construct. void emitUserDefinedMapper(const OMPDeclareMapperDecl *D, CodeGenFunction *CGF = nullptr); + /// Get the function for the specified user-defined mapper, if any. + virtual llvm::Function * + getUserDefinedMapperFunc(const OMPDeclareMapperDecl *D); /// Emits outlined function for the specified OpenMP parallel directive /// \a D. This outlined function has type void(*)(kmp_int32 *ThreadID, @@ -1510,6 +1513,8 @@ llvm::Value *SizesArray = nullptr; /// The array of map types passed to the runtime library. llvm::Value *MapTypesArray = nullptr; + /// The array of user-defined mappers passed to the runtime library. + llvm::Value *MappersArray = nullptr; /// The total number of pointers passed to the runtime library. unsigned NumberOfPtrs = 0u; /// Map between the a declaration of a capture and the corresponding base @@ -1525,12 +1530,13 @@ PointersArray = nullptr; SizesArray = nullptr; MapTypesArray = nullptr; + MappersArray = nullptr; NumberOfPtrs = 0u; } /// Return true if the current target data information has valid arrays. bool isValid() { return BasePointersArray && PointersArray && SizesArray && - MapTypesArray && NumberOfPtrs; + MapTypesArray && MappersArray && NumberOfPtrs; } bool requiresDevicePointerInfo() { return RequiresDevicePointerInfo; } }; Index: lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenMPRuntime.cpp +++ lib/CodeGen/CGOpenMPRuntime.cpp @@ -27,6 +27,7 @@ #include "llvm/Support/Format.h" #include "llvm/Support/raw_ostream.h" #include +#include using namespace clang; using namespace CodeGen; @@ -739,7 +740,7 @@ // *arg_types); OMPRTL__tgt_target_data_begin_nowait, // Call to void __tgt_target_data_end(int64_t device_id, int32_t arg_num, - // void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types); + // void** args_base, void **args, int64_t *arg_sizes, int64_t *arg_types); OMPRTL__tgt_target_data_end, // Call to void __tgt_target_data_end_nowait(int64_t device_id, int32_t // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t @@ -752,6 +753,48 @@ // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t // *arg_types); OMPRTL__tgt_target_data_update_nowait, + // Call to int32_t __tgt_target_mapper(int64_t device_id, void *host_ptr, + // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + OMPRTL__tgt_target_mapper, + // Call to int32_t __tgt_target_nowait_mapper(int64_t device_id, void + // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t + // *arg_sizes, int64_t *arg_types, void **arg_mappers); + OMPRTL__tgt_target_nowait_mapper, + // Call to int32_t __tgt_target_teams_mapper(int64_t device_id, void + // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t + // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams, + // int32_t thread_limit); + OMPRTL__tgt_target_teams_mapper, + // Call to int32_t __tgt_target_teams_nowait_mapper(int64_t device_id, void + // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t + // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams, + // int32_t thread_limit); + OMPRTL__tgt_target_teams_nowait_mapper, + // Call to void __tgt_target_data_begin_mapper(int64_t device_id, int32_t + // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + OMPRTL__tgt_target_data_begin_mapper, + // Call to void __tgt_target_data_begin_nowait_mapper(int64_t device_id, + // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + OMPRTL__tgt_target_data_begin_nowait_mapper, + // Call to void __tgt_target_data_end_mapper(int64_t device_id, int32_t + // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + OMPRTL__tgt_target_data_end_mapper, + // Call to void __tgt_target_data_end_nowait_mapper(int64_t device_id, int32_t + // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + OMPRTL__tgt_target_data_end_nowait_mapper, + // Call to void __tgt_target_data_update_mapper(int64_t device_id, int32_t + // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + OMPRTL__tgt_target_data_update_mapper, + // Call to void __tgt_target_data_update_nowait_mapper(int64_t device_id, + // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + OMPRTL__tgt_target_data_update_nowait_mapper, // 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 @@ -2470,6 +2513,179 @@ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_nowait"); break; } + case OMPRTL__tgt_target_mapper: { + // Build int32_t __tgt_target_mapper(int64_t device_id, void *host_ptr, + // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, + // int64_t *arg_types, void **arg_mappers); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.VoidPtrTy, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_mapper"); + break; + } + case OMPRTL__tgt_target_nowait_mapper: { + // Build int32_t __tgt_target_nowait_mapper(int64_t device_id, void + // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t + // *arg_sizes, int64_t *arg_types, void **arg_mappers); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.VoidPtrTy, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_nowait_mapper"); + break; + } + case OMPRTL__tgt_target_teams_mapper: { + // Build int32_t __tgt_target_teams_mapper(int64_t device_id, void + // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t + // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams, + // int32_t thread_limit); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.VoidPtrTy, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy, + CGM.Int32Ty, + CGM.Int32Ty}; + auto *FnTy = + llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_teams_mapper"); + break; + } + case OMPRTL__tgt_target_teams_nowait_mapper: { + // Build int32_t __tgt_target_teams_nowait_mapper(int64_t device_id, void + // *host_ptr, int32_t arg_num, void** args_base, void **args, int64_t + // *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t num_teams, + // int32_t thread_limit); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.VoidPtrTy, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy, + CGM.Int32Ty, + CGM.Int32Ty}; + auto *FnTy = + llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_teams_nowait_mapper"); + break; + } + case OMPRTL__tgt_target_data_begin_mapper: { + // Build void __tgt_target_data_begin_mapper(int64_t device_id, int32_t + // arg_num, void **args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_begin_mapper"); + break; + } + case OMPRTL__tgt_target_data_begin_nowait_mapper: { + // Build void __tgt_target_data_begin_nowait_mapper(int64_t device_id, + // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, + // int64_t *arg_types, void **arg_mappers); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, + "__tgt_target_data_begin_nowait_mapper"); + break; + } + case OMPRTL__tgt_target_data_end_mapper: { + // Build void __tgt_target_data_end_mapper(int64_t device_id, int32_t + // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_end_mapper"); + break; + } + case OMPRTL__tgt_target_data_end_nowait_mapper: { + // Build void __tgt_target_data_end_nowait_mapper(int64_t device_id, int32_t + // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false); + RTLFn = + CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_end_nowait_mapper"); + break; + } + case OMPRTL__tgt_target_data_update_mapper: { + // Build void __tgt_target_data_update_mapper(int64_t device_id, int32_t + // arg_num, void** args_base, void **args, int64_t *arg_sizes, int64_t + // *arg_types, void **arg_mappers); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_mapper"); + break; + } + case OMPRTL__tgt_target_data_update_nowait_mapper: { + // Build void __tgt_target_data_update_nowait_mapper(int64_t device_id, + // int32_t arg_num, void** args_base, void **args, int64_t *arg_sizes, + // int64_t *arg_types, void **arg_mappers); + llvm::Type *TypeParams[] = {CGM.Int64Ty, + CGM.Int32Ty, + CGM.VoidPtrPtrTy, + CGM.VoidPtrPtrTy, + CGM.Int64Ty->getPointerTo(), + CGM.Int64Ty->getPointerTo(), + CGM.VoidPtrPtrTy}; + auto *FnTy = + llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false); + RTLFn = CGM.CreateRuntimeFunction(FnTy, + "__tgt_target_data_update_nowait_mapper"); + break; + } case OMPRTL__tgt_mapper_num_components: { // Build int64_t __tgt_mapper_num_components(void *rt_mapper_handle); llvm::Type *TypeParams[] = {CGM.VoidPtrTy}; @@ -7154,6 +7370,7 @@ using MapBaseValuesArrayTy = SmallVector; using MapValuesArrayTy = SmallVector; using MapFlagsArrayTy = SmallVector; + using MapMappersArrayTy = SmallVector; /// Map between a struct and the its lowest & highest elements which have been /// mapped. @@ -7175,15 +7392,17 @@ ArrayRef MapModifiers; bool ReturnDevicePointer = false; bool IsImplicit = false; + const ValueDecl *Mapper = nullptr; MapInfo() = default; MapInfo( OMPClauseMappableExprCommon::MappableExprComponentListRef Components, OpenMPMapClauseKind MapType, - ArrayRef MapModifiers, - bool ReturnDevicePointer, bool IsImplicit) + ArrayRef MapModifiers, bool ReturnDevicePointer, + bool IsImplicit, const ValueDecl *Mapper = nullptr) : Components(Components), MapType(MapType), MapModifiers(MapModifiers), - ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit) {} + ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit), + Mapper(Mapper) {} }; /// If use_device_ptr is used on a pointer which is a struct member and there @@ -7218,7 +7437,7 @@ SmallVector> DevPointersMap; - llvm::Value *getExprTypeSize(const Expr *E) const { + llvm::Value *getExprTypeSize(const Expr *E, bool hasMapper) const { QualType ExprTy = E->getType().getCanonicalType(); // Reference types are ignored for mapping purposes. @@ -7235,8 +7454,14 @@ // If there is no length associated with the expression, that means we // are using the whole length of the base. - if (!OAE->getLength() && OAE->getColonLoc().isValid()) - return CGF.getTypeSize(BaseTy); + if (!OAE->getLength() && OAE->getColonLoc().isValid()) { + // In case that a user-defined mapper is attached, its size is the + // number of array elements instead of the number of total bytes. + if (hasMapper) + return CGF.Builder.getInt64(1); + else + return CGF.getTypeSize(BaseTy); + } llvm::Value *ElemSize; if (const auto *PTy = BaseTy->getAs()) { @@ -7249,15 +7474,31 @@ // If we don't have a length at this point, that is because we have an // array section with a single element. - if (!OAE->getLength()) - return ElemSize; + if (!OAE->getLength()) { + // In case that a user-defined mapper is attached, its size is the + // number of array elements instead of the number of total bytes. + if (hasMapper) + return CGF.Builder.getInt64(1); + else + return ElemSize; + } llvm::Value *LengthVal = CGF.EmitScalarExpr(OAE->getLength()); LengthVal = CGF.Builder.CreateIntCast(LengthVal, CGF.SizeTy, /*isSigned=*/false); - return CGF.Builder.CreateNUWMul(LengthVal, ElemSize); + // In case that a user-defined mapper is attached, its size is the + // number of array elements instead of the number of total bytes. + if (hasMapper) + return LengthVal; + else + return CGF.Builder.CreateNUWMul(LengthVal, ElemSize); } - return CGF.getTypeSize(ExprTy); + // In case that a user-defined mapper is attached, its size is the + // number of array elements instead of the number of total bytes. + if (hasMapper) + return CGF.Builder.getInt64(1); + else + return CGF.getTypeSize(ExprTy); } /// Return the corresponding bits for a given map clause modifier. Add @@ -7344,18 +7585,18 @@ return ConstLength.getSExtValue() != 1; } - /// Generate the base pointers, section pointers, sizes and map type - /// bits for the provided map type, map modifier, and expression components. - /// \a IsFirstComponent should be set to true if the provided set of - /// components is the first associated with a capture. + /// Generate the base pointers, section pointers, sizes, map type bits, and + /// mappers for the provided map type, map modifier, and expression + /// components. \a IsFirstComponent should be set to true if the provided set + /// of components is the first associated with a capture. void generateInfoForComponentList( - OpenMPMapClauseKind MapType, - ArrayRef MapModifiers, + OpenMPMapClauseKind MapType, ArrayRef MapModifiers, OMPClauseMappableExprCommon::MappableExprComponentListRef Components, MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, - StructRangeInfoTy &PartialStruct, bool IsFirstComponentList, - bool IsImplicit, + MapMappersArrayTy &Mappers, StructRangeInfoTy &PartialStruct, + bool IsFirstComponentList, bool IsImplicit, + const ValueDecl *Mapper = nullptr, ArrayRef OverlappedElements = llvm::None) const { // The following summarizes what has to be generated for each map and the @@ -7699,6 +7940,7 @@ Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); Types.push_back(Flags); + Mappers.push_back(nullptr); LB = CGF.Builder.CreateConstGEP(ComponentLB, 1); } BasePointers.push_back(BP.getPointer()); @@ -7710,14 +7952,29 @@ Sizes.push_back( CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); Types.push_back(Flags); + Mappers.push_back(nullptr); break; } - llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression()); if (!IsMemberPointer) { BasePointers.push_back(BP.getPointer()); Pointers.push_back(LB.getPointer()); - Sizes.push_back( - CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true)); + + // If Mapper is valid, the last component inherits the mapper. + bool hasMapper = Mapper && Next == CE; + llvm::Value *Size = + getExprTypeSize(I->getAssociatedExpression(), hasMapper); + Sizes.push_back(CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, + /*isSigned=*/true)); + if (hasMapper) + Mappers.push_back(Mapper); + else + Mappers.push_back(nullptr); + std::cerr << "HH " << "\n"; + if (hasMapper) { + std::cerr << "AM: " << Mappers.size() << " "; + I->getAssociatedExpression()->dump(); + //Size->dump(); + } // We need to add a pointer flag for each map that comes from the // same expression except for the first one. We also need to signal @@ -7898,7 +8155,7 @@ // Extract device pointer clause information. for (const auto *C : Dir.getClausesOfKind()) for (auto L : C->component_lists()) - DevPointersMap[L.first].push_back(L.second); + DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L)); } /// Constructor for the declare mapper directive. @@ -7910,13 +8167,16 @@ /// individual struct members. void emitCombinedEntry(MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, - MapFlagsArrayTy &Types, MapFlagsArrayTy &CurTypes, + MapFlagsArrayTy &Types, MapMappersArrayTy &Mappers, + MapFlagsArrayTy &CurTypes, const StructRangeInfoTy &PartialStruct) const { // Base is the base of the struct BasePointers.push_back(PartialStruct.Base.getPointer()); // Pointer is the address of the lowest element llvm::Value *LB = PartialStruct.LowestElem.second.getPointer(); Pointers.push_back(LB); + // There should not be a mapper for a combined entry. + Mappers.push_back(nullptr); // Size is (addr of {highest+1} element) - (addr of lowest element) llvm::Value *HB = PartialStruct.HighestElem.second.getPointer(); llvm::Value *HAddr = CGF.Builder.CreateConstGEP1_32(HB, /*Idx0=*/1); @@ -7940,13 +8200,14 @@ setCorrectMemberOfFlag(M, MemberOfFlag); } - /// Generate all the base pointers, section pointers, sizes and map - /// types for the extracted mappable expressions. Also, for each item that + /// Generate all the base pointers, section pointers, sizes, map types, and + /// mappers for the extracted mappable expressions. Also, for each item that /// relates with a device pointer, a pair of the relevant declaration and /// index where it occurs is appended to the device pointers info array. void generateAllInfo(MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, - MapFlagsArrayTy &Types) const { + MapFlagsArrayTy &Types, + MapMappersArrayTy &Mappers) const { // 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. @@ -7954,35 +8215,37 @@ // 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); - }; + auto &&InfoGen = + [&Info](const ValueDecl *D, + OMPClauseMappableExprCommon::MappableExprComponentListRef L, + OpenMPMapClauseKind MapType, + ArrayRef MapModifiers, + bool ReturnDevicePointer, bool IsImplicit, + const ValueDecl *Mapper) { + const ValueDecl *VD = + D ? cast(D->getCanonicalDecl()) : nullptr; + Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer, + IsImplicit, Mapper); + }; assert(CurDir.is() && "Expect a executable directive"); const auto *CurExecDir = CurDir.get(); for (const auto *C : CurExecDir->getClausesOfKind()) for (const auto &L : C->component_lists()) { - InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifiers(), - /*ReturnDevicePointer=*/false, C->isImplicit()); + InfoGen(std::get<0>(L), std::get<1>(L), C->getMapType(), + C->getMapTypeModifiers(), /*ReturnDevicePointer=*/false, + C->isImplicit(), std::get<2>(L)); } for (const auto *C : CurExecDir->getClausesOfKind()) for (const auto &L : C->component_lists()) { - InfoGen(L.first, L.second, OMPC_MAP_to, llvm::None, - /*ReturnDevicePointer=*/false, C->isImplicit()); + InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_to, llvm::None, + /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L)); } for (const auto *C : CurExecDir->getClausesOfKind()) for (const auto &L : C->component_lists()) { - InfoGen(L.first, L.second, OMPC_MAP_from, llvm::None, - /*ReturnDevicePointer=*/false, C->isImplicit()); + InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_from, llvm::None, + /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L)); } // Look at the use_device_ptr clause information and mark the existing map @@ -7997,10 +8260,13 @@ for (const auto *C : CurExecDir->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(); + OMPClauseMappableExprCommon::MappableExprComponentListRef Components = + std::get<1>(L); + assert(!Components.empty() && + "Not expecting empty list of components!"); + const ValueDecl *VD = Components.back().getAssociatedDeclaration(); VD = cast(VD->getCanonicalDecl()); - const Expr *IE = L.second.back().getAssociatedExpression(); + const Expr *IE = Components.back().getAssociatedExpression(); // If the first component is a member expression, we have to look into // 'this', which maps to null in the map of map information. Otherwise // look directly for the information. @@ -8032,8 +8298,8 @@ // Nonetheless, generateInfoForComponentList must be called to take // the pointer into account for the calculation of the range of the // partial struct. - InfoGen(nullptr, L.second, OMPC_MAP_unknown, llvm::None, - /*ReturnDevicePointer=*/false, C->isImplicit()); + InfoGen(nullptr, Components, OMPC_MAP_unknown, llvm::None, + /*ReturnDevicePointer=*/false, C->isImplicit(), nullptr); DeferredInfo[nullptr].emplace_back(IE, VD); } else { llvm::Value *Ptr = @@ -8042,6 +8308,7 @@ Pointers.push_back(Ptr); Sizes.push_back(llvm::Constant::getNullValue(CGF.Int64Ty)); Types.push_back(OMP_MAP_RETURN_PARAM | OMP_MAP_TARGET_PARAM); + Mappers.push_back(nullptr); } } } @@ -8056,6 +8323,7 @@ MapValuesArrayTy CurPointers; MapValuesArrayTy CurSizes; MapFlagsArrayTy CurTypes; + MapMappersArrayTy CurMappers; StructRangeInfoTy PartialStruct; for (const MapInfo &L : M.second) { @@ -8064,10 +8332,10 @@ // 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); + generateInfoForComponentList( + L.MapType, L.MapModifiers, L.Components, CurBasePointers, + CurPointers, CurSizes, CurTypes, CurMappers, PartialStruct, + IsFirstComponentList, L.IsImplicit, L.Mapper); // If this entry relates with a device pointer, set the relevant // declaration and add the 'return pointer' flag. @@ -8102,29 +8370,31 @@ // correct value of MEMBER_OF. CurTypes.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_RETURN_PARAM | OMP_MAP_MEMBER_OF); + CurMappers.push_back(nullptr); } } // 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); + emitCombinedEntry(BasePointers, Pointers, Sizes, Types, Mappers, + 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()); + Mappers.append(CurMappers.begin(), CurMappers.end()); } } - /// Generate all the base pointers, section pointers, sizes and map types for - /// the extracted map clauses of user-defined mapper. + /// Generate all the base pointers, section pointers, sizes, map types, and + /// mappers for the extracted map clauses of user-defined mapper. void generateAllInfoForMapper(MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, - MapValuesArrayTy &Sizes, - MapFlagsArrayTy &Types) const { + MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, + MapMappersArrayTy &Mappers) const { assert(CurDir.is() && "Expect a declare mapper directive"); const auto *CurMapperDir = CurDir.get(); @@ -8133,25 +8403,17 @@ // 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); - }; - + // Fill the information map for map clauses. for (const auto *C : CurMapperDir->clauselists()) { - const auto *MC = cast(C); + 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()); + const ValueDecl *VD = + std::get<0>(L) ? cast(std::get<0>(L)->getCanonicalDecl()) + : nullptr; + // Get the corresponding user-defined mapper. + Info[VD].emplace_back( + std::get<1>(L), MC->getMapType(), MC->getMapTypeModifiers(), + /*ReturnDevicePointer=*/false, MC->isImplicit(), std::get<2>(L)); } } @@ -8165,29 +8427,31 @@ MapValuesArrayTy CurPointers; MapValuesArrayTy CurSizes; MapFlagsArrayTy CurTypes; + MapMappersArrayTy CurMappers; 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); + generateInfoForComponentList( + L.MapType, L.MapModifiers, L.Components, CurBasePointers, + CurPointers, CurSizes, CurTypes, CurMappers, PartialStruct, + IsFirstComponentList, L.IsImplicit, L.Mapper); 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); + emitCombinedEntry(BasePointers, Pointers, Sizes, Types, Mappers, + 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()); + Mappers.append(CurMappers.begin(), CurMappers.end()); } } @@ -8195,7 +8459,7 @@ void generateInfoForLambdaCaptures( const ValueDecl *VD, llvm::Value *Arg, MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, - MapFlagsArrayTy &Types, + MapFlagsArrayTy &Types, MapMappersArrayTy &Mappers, llvm::DenseMap &LambdaPointers) const { const auto *RD = VD->getType() .getCanonicalType() @@ -8221,6 +8485,7 @@ CGF.Int64Ty, /*isSigned=*/true)); Types.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_LITERAL | OMP_MAP_MEMBER_OF | OMP_MAP_IMPLICIT); + Mappers.push_back(nullptr); } for (const LambdaCapture &LC : RD->captures()) { if (!LC.capturesVariable()) @@ -8249,6 +8514,7 @@ } Types.push_back(OMP_MAP_PTR_AND_OBJ | OMP_MAP_LITERAL | OMP_MAP_MEMBER_OF | OMP_MAP_IMPLICIT); + Mappers.push_back(nullptr); } } @@ -8281,13 +8547,14 @@ } } - /// Generate the base pointers, section pointers, sizes and map types - /// associated to a given capture. + /// Generate the base pointers, section pointers, sizes, map types, and + /// mappers associated to a given capture. void generateInfoForCapture(const CapturedStmt::Capture *Cap, llvm::Value *Arg, MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, + MapMappersArrayTy &Mappers, StructRangeInfoTy &PartialStruct) const { assert(!Cap->capturesVariableArrayType() && "Not expecting to generate map info for a variable array type!"); @@ -8307,25 +8574,30 @@ CGF.Builder.CreateIntCast(CGF.getTypeSize(CGF.getContext().VoidPtrTy), CGF.Int64Ty, /*isSigned=*/true)); Types.push_back(OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM); + Mappers.push_back(nullptr); return; } using MapData = std::tuple, bool>; + OpenMPMapClauseKind, ArrayRef, bool, + const ValueDecl *>; SmallVector DeclComponentLists; assert(CurDir.is() && "Expect a executable directive"); const auto *CurExecDir = CurDir.get(); + std::cerr << "C " << "\n"; for (const auto *C : CurExecDir->getClausesOfKind()) { for (const auto &L : C->decl_component_lists(VD)) { - assert(L.first == VD && - "We got information for the wrong declaration??"); - assert(!L.second.empty() && + const ValueDecl *VDecl, *Mapper; + OMPClauseMappableExprCommon::MappableExprComponentListRef Components; + std::tie(VDecl, Components, Mapper) = L; + assert(VDecl == VD && "We got information for the wrong declaration??"); + assert(!Components.empty() && "Not expecting declaration with no component lists."); - DeclComponentLists.emplace_back(L.second, C->getMapType(), + DeclComponentLists.emplace_back(Components, C->getMapType(), C->getMapTypeModifiers(), - C->isImplicit()); + C->isImplicit(), Mapper); } } @@ -8342,11 +8614,12 @@ OpenMPMapClauseKind MapType; ArrayRef MapModifiers; bool IsImplicit; - std::tie(Components, MapType, MapModifiers, IsImplicit) = L; + const ValueDecl *Mapper; + std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L; ++Count; for (const MapData &L1 : makeArrayRef(DeclComponentLists).slice(Count)) { OMPClauseMappableExprCommon::MappableExprComponentListRef Components1; - std::tie(Components1, MapType, MapModifiers, IsImplicit) = L1; + std::tie(Components1, MapType, MapModifiers, IsImplicit, Mapper) = L1; auto CI = Components.rbegin(); auto CE = Components.rend(); auto SI = Components1.rbegin(); @@ -8432,14 +8705,15 @@ OpenMPMapClauseKind MapType; ArrayRef MapModifiers; bool IsImplicit; - std::tie(Components, MapType, MapModifiers, IsImplicit) = L; + const ValueDecl *Mapper; + std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L; ArrayRef OverlappedComponents = Pair.getSecond(); bool IsFirstComponentList = true; generateInfoForComponentList(MapType, MapModifiers, Components, BasePointers, Pointers, Sizes, Types, - PartialStruct, IsFirstComponentList, - IsImplicit, OverlappedComponents); + Mappers, PartialStruct, IsFirstComponentList, + IsImplicit, Mapper, OverlappedComponents); } // Go through other elements without overlapped elements. bool IsFirstComponentList = OverlappedData.empty(); @@ -8448,23 +8722,25 @@ OpenMPMapClauseKind MapType; ArrayRef MapModifiers; bool IsImplicit; - std::tie(Components, MapType, MapModifiers, IsImplicit) = L; + const ValueDecl *Mapper; + std::tie(Components, MapType, MapModifiers, IsImplicit, Mapper) = L; auto It = OverlappedData.find(&L); if (It == OverlappedData.end()) generateInfoForComponentList(MapType, MapModifiers, Components, BasePointers, Pointers, Sizes, Types, - PartialStruct, IsFirstComponentList, - IsImplicit); + Mappers, PartialStruct, + IsFirstComponentList, IsImplicit, Mapper); IsFirstComponentList = false; } } - /// Generate the base pointers, section pointers, sizes and map types - /// associated with the declare target link variables. + /// Generate the base pointers, section pointers, sizes, map types, and + /// mappers associated with the declare target link variables. void generateInfoForDeclareTargetLink(MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, - MapFlagsArrayTy &Types) const { + MapFlagsArrayTy &Types, + MapMappersArrayTy &Mappers) const { assert(CurDir.is() && "Expect a executable directive"); const auto *CurExecDir = CurDir.get(); @@ -8472,9 +8748,9 @@ // but "declare target link" global variables. for (const auto *C : CurExecDir->getClausesOfKind()) { for (const auto &L : C->component_lists()) { - if (!L.first) + if (!std::get<0>(L)) continue; - const auto *VD = dyn_cast(L.first); + const auto *VD = dyn_cast(std::get<0>(L)); if (!VD) continue; llvm::Optional Res = @@ -8484,8 +8760,8 @@ continue; StructRangeInfoTy PartialStruct; generateInfoForComponentList( - C->getMapType(), C->getMapTypeModifiers(), L.second, BasePointers, - Pointers, Sizes, Types, PartialStruct, + C->getMapType(), C->getMapTypeModifiers(), std::get<1>(L), + BasePointers, Pointers, Sizes, Types, Mappers, PartialStruct, /*IsFirstComponentList=*/true, C->isImplicit()); assert(!PartialStruct.Base.isValid() && "No partial structs for declare target link expected."); @@ -8500,7 +8776,8 @@ MapBaseValuesArrayTy &CurBasePointers, MapValuesArrayTy &CurPointers, MapValuesArrayTy &CurSizes, - MapFlagsArrayTy &CurMapTypes) const { + MapFlagsArrayTy &CurMapTypes, + MapMappersArrayTy &CurMappers) const { bool IsImplicit = true; // Do the default mapping. if (CI.capturesThis()) { @@ -8575,6 +8852,9 @@ // Add flag stating this is an implicit map. if (IsImplicit) CurMapTypes.back() |= OMP_MAP_IMPLICIT; + + // No user-defined mapper for default mapping. + CurMappers.push_back(nullptr); } }; } // anonymous namespace @@ -8588,6 +8868,7 @@ MappableExprsHandler::MapValuesArrayTy &Pointers, MappableExprsHandler::MapValuesArrayTy &Sizes, MappableExprsHandler::MapFlagsArrayTy &MapTypes, + MappableExprsHandler::MapMappersArrayTy &Mappers, CGOpenMPRuntime::TargetDataInfo &Info) { CodeGenModule &CGM = CGF.CGM; ASTContext &Ctx = CGF.getContext(); @@ -8615,6 +8896,8 @@ CGF.CreateMemTemp(PointerArrayType, ".offload_baseptrs").getPointer(); Info.PointersArray = CGF.CreateMemTemp(PointerArrayType, ".offload_ptrs").getPointer(); + Info.MappersArray = + CGF.CreateMemTemp(PointerArrayType, ".offload_mappers").getPointer(); // If we don't have any VLA types or other types that require runtime // evaluation, we can use a constant array for the map sizes, otherwise we @@ -8694,16 +8977,30 @@ CGF.Builder.CreateIntCast(Sizes[I], CGM.Int64Ty, /*isSigned=*/true), SAddr); } + + // Fill up the mapper array. + llvm::Value *MFunc = llvm::ConstantPointerNull::get(CGM.VoidPtrTy); + if (Mappers[I]) + MFunc = CGM.getOpenMPRuntime().getUserDefinedMapperFunc( + cast(Mappers[I])); + llvm::Value *M = CGF.Builder.CreateConstInBoundsGEP2_32( + llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs), + Info.MappersArray, 0, I); + M = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast( + M, MFunc->getType()->getPointerTo(/*AddrSpace=*/0)); + Address MAddr(M, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy)); + CGF.Builder.CreateStore(MFunc, MAddr); } } } /// Emit the arguments to be passed to the runtime library based on the -/// arrays of pointers, sizes and map types. +/// arrays of base pointers, pointers, sizes, map types, and mappers. static void emitOffloadingArraysArgument( CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg, llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg, - llvm::Value *&MapTypesArrayArg, CGOpenMPRuntime::TargetDataInfo &Info) { + llvm::Value *&MapTypesArrayArg, llvm::Value *&MappersArrayArg, + CGOpenMPRuntime::TargetDataInfo &Info) { CodeGenModule &CGM = CGF.CGM; if (Info.NumberOfPtrs) { BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32( @@ -8723,12 +9020,17 @@ Info.MapTypesArray, /*Idx0=*/0, /*Idx1=*/0); + MappersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32( + llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs), + Info.MappersArray, + /*Idx0=*/0, /*Idx1=*/0); } else { BasePointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy); PointersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy); SizesArrayArg = llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo()); MapTypesArrayArg = llvm::ConstantPointerNull::get(CGM.Int64Ty->getPointerTo()); + MappersArrayArg = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy); } } @@ -8942,6 +9244,7 @@ // Emit the loop body block. MapperCGF.EmitBlock(BodyBB); + llvm::BasicBlock *LastBB = BodyBB; llvm::PHINode *PtrPHI = MapperCGF.Builder.CreatePHI( PtrBegin->getType(), 2, "omp.arraymap.ptrcurrent"); PtrPHI->addIncoming(PtrBegin, EntryBB); @@ -8963,8 +9266,10 @@ MappableExprsHandler::MapValuesArrayTy Pointers; MappableExprsHandler::MapValuesArrayTy Sizes; MappableExprsHandler::MapFlagsArrayTy MapTypes; + MappableExprsHandler::MapMappersArrayTy Mappers; MappableExprsHandler MEHandler(*D, MapperCGF); - MEHandler.generateAllInfoForMapper(BasePointers, Pointers, Sizes, MapTypes); + MEHandler.generateAllInfoForMapper(BasePointers, Pointers, Sizes, MapTypes, + Mappers); // Call the runtime API __tgt_mapper_num_components to get the number of // pre-existing components. @@ -9062,6 +9367,7 @@ MapperCGF.Builder.getInt64(~MappableExprsHandler::OMP_MAP_TO)); // In case of tofrom, do nothing. MapperCGF.EmitBlock(EndBB); + LastBB = EndBB; llvm::PHINode *CurMapType = MapperCGF.Builder.CreatePHI(CGM.Int64Ty, 4, "omp.maptype"); CurMapType->addIncoming(AllocMapType, AllocBB); @@ -9069,22 +9375,28 @@ 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); + if (Mappers[I]) { + // Call the corresponding mapper function. + llvm::Function *MapperFunc = + getUserDefinedMapperFunc(cast(Mappers[I])); + assert(MapperFunc && "Expect a valid mapper function is available."); + MapperCGF.Builder.CreateCall(MapperFunc, OffloadingArgs); + } else { + // Call the runtime API __tgt_push_mapper_component to fill up the runtime + // data structure. + 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); + PtrPHI->addIncoming(PtrNext, LastBB); llvm::Value *IsDone = MapperCGF.Builder.CreateICmpEQ(PtrNext, PtrEnd, "omp.arraymap.isdone"); llvm::BasicBlock *ExitBB = MapperCGF.createBasicBlock("omp.arraymap.exit"); @@ -9159,6 +9471,15 @@ createRuntimeFunction(OMPRTL__tgt_push_mapper_component), OffloadingArgs); } +llvm::Function * +CGOpenMPRuntime::getUserDefinedMapperFunc(const OMPDeclareMapperDecl *D) { + auto I = UDMMap.find(D); + if (I != UDMMap.end()) + return I->second; + emitUserDefinedMapper(D); + return UDMMap.lookup(D); +} + void CGOpenMPRuntime::emitTargetNumIterationsCall( CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *Device, const llvm::function_refdump(); // VLA sizes are passed to the outlined region by copy and do not have map // information associated. @@ -9374,20 +9701,23 @@ CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_LITERAL | MappableExprsHandler::OMP_MAP_TARGET_PARAM | MappableExprsHandler::OMP_MAP_IMPLICIT); + CurMappers.push_back(nullptr); } else { // If we have any information in the map clause, we use it, otherwise we // just do a default mapping. MEHandler.generateInfoForCapture(CI, *CV, CurBasePointers, CurPointers, - CurSizes, CurMapTypes, PartialStruct); + CurSizes, CurMapTypes, CurMappers, + PartialStruct); if (CurBasePointers.empty()) MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers, - CurPointers, CurSizes, CurMapTypes); + CurPointers, CurSizes, CurMapTypes, + CurMappers); // Generate correct mapping for variables captured by reference in // lambdas. if (CI->capturesVariable()) MEHandler.generateInfoForLambdaCaptures( CI->getCapturedVar(), *CV, CurBasePointers, CurPointers, CurSizes, - CurMapTypes, LambdaPointers); + CurMapTypes, CurMappers, LambdaPointers); } // We expect to have at least an element of information for this capture. assert(!CurBasePointers.empty() && @@ -9395,19 +9725,21 @@ assert(CurBasePointers.size() == CurPointers.size() && CurBasePointers.size() == CurSizes.size() && CurBasePointers.size() == CurMapTypes.size() && + CurBasePointers.size() == CurMappers.size() && "Inconsistent map information sizes!"); // 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()) MEHandler.emitCombinedEntry(BasePointers, Pointers, Sizes, MapTypes, - CurMapTypes, PartialStruct); + Mappers, CurMapTypes, 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()); MapTypes.append(CurMapTypes.begin(), CurMapTypes.end()); + Mappers.append(CurMappers.begin(), CurMappers.end()); } // Adjust MEMBER_OF flags for the lambdas captures. MEHandler.adjustMemberOfForLambdaCaptures(LambdaPointers, BasePointers, @@ -9415,20 +9747,22 @@ // Map other list items in the map clause which are not captured variables // but "declare target link" global variables. MEHandler.generateInfoForDeclareTargetLink(BasePointers, Pointers, Sizes, - MapTypes); + MapTypes, Mappers); TargetDataInfo Info; // Fill up the arrays and create the arguments. - emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info); + emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Mappers, + Info); emitOffloadingArraysArgument(CGF, Info.BasePointersArray, Info.PointersArray, Info.SizesArray, - Info.MapTypesArray, Info); + Info.MapTypesArray, Info.MappersArray, Info); InputInfo.NumberOfTargetItems = Info.NumberOfPtrs; InputInfo.BasePointersArray = Address(Info.BasePointersArray, CGM.getPointerAlign()); InputInfo.PointersArray = Address(Info.PointersArray, CGM.getPointerAlign()); InputInfo.SizesArray = Address(Info.SizesArray, CGM.getPointerAlign()); + InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign()); MapTypesArray = Info.MapTypesArray; if (RequiresOuterTask) CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo); @@ -9983,20 +10317,24 @@ MappableExprsHandler::MapValuesArrayTy Pointers; MappableExprsHandler::MapValuesArrayTy Sizes; MappableExprsHandler::MapFlagsArrayTy MapTypes; + MappableExprsHandler::MapMappersArrayTy Mappers; // Get map clause information. - MappableExprsHandler MCHandler(D, CGF); - MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes); + MappableExprsHandler MEHandler(D, CGF); + MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Mappers); // Fill up the arrays and create the arguments. - emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info); + emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Mappers, + Info); llvm::Value *BasePointersArrayArg = nullptr; llvm::Value *PointersArrayArg = nullptr; llvm::Value *SizesArrayArg = nullptr; llvm::Value *MapTypesArrayArg = nullptr; + llvm::Value *MappersArrayArg = nullptr; emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg, - SizesArrayArg, MapTypesArrayArg, Info); + SizesArrayArg, MapTypesArrayArg, + MappersArrayArg, Info); // Emit device ID if any. llvm::Value *DeviceID = nullptr; @@ -10011,10 +10349,11 @@ llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs); llvm::Value *OffloadingArgs[] = { - DeviceID, PointerNum, BasePointersArrayArg, - PointersArrayArg, SizesArrayArg, MapTypesArrayArg}; - CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_target_data_begin), - OffloadingArgs); + DeviceID, PointerNum, BasePointersArrayArg, PointersArrayArg, + SizesArrayArg, MapTypesArrayArg, MappersArrayArg}; + CGF.EmitRuntimeCall( + createRuntimeFunction(OMPRTL__tgt_target_data_begin_mapper), + OffloadingArgs); // If device pointer privatization is required, emit the body of the region // here. It will have to be duplicated: with and without privatization. @@ -10031,8 +10370,10 @@ llvm::Value *PointersArrayArg = nullptr; llvm::Value *SizesArrayArg = nullptr; llvm::Value *MapTypesArrayArg = nullptr; + llvm::Value *MappersArrayArg = nullptr; emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg, - SizesArrayArg, MapTypesArrayArg, Info); + SizesArrayArg, MapTypesArrayArg, + MappersArrayArg, Info); // Emit device ID if any. llvm::Value *DeviceID = nullptr; @@ -10047,10 +10388,11 @@ llvm::Value *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs); llvm::Value *OffloadingArgs[] = { - DeviceID, PointerNum, BasePointersArrayArg, - PointersArrayArg, SizesArrayArg, MapTypesArrayArg}; - CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_target_data_end), - OffloadingArgs); + DeviceID, PointerNum, BasePointersArrayArg, PointersArrayArg, + SizesArrayArg, MapTypesArrayArg, MappersArrayArg}; + CGF.EmitRuntimeCall( + createRuntimeFunction(OMPRTL__tgt_target_data_end_mapper), + OffloadingArgs); }; // If we need device pointer privatization, we need to emit the body of the @@ -10124,24 +10466,25 @@ InputInfo.BasePointersArray.getPointer(), InputInfo.PointersArray.getPointer(), InputInfo.SizesArray.getPointer(), - MapTypesArray}; + MapTypesArray, + InputInfo.MappersArray.getPointer()}; - // Select the right runtime function call for each expected standalone + // Select the right runtime function call for each standalone // directive. const bool HasNowait = D.hasClausesOfKind(); OpenMPRTLFunction RTLFn; switch (D.getDirectiveKind()) { case OMPD_target_enter_data: - RTLFn = HasNowait ? OMPRTL__tgt_target_data_begin_nowait - : OMPRTL__tgt_target_data_begin; + RTLFn = HasNowait ? OMPRTL__tgt_target_data_begin_nowait_mapper + : OMPRTL__tgt_target_data_begin_mapper; break; case OMPD_target_exit_data: - RTLFn = HasNowait ? OMPRTL__tgt_target_data_end_nowait - : OMPRTL__tgt_target_data_end; + RTLFn = HasNowait ? OMPRTL__tgt_target_data_end_nowait_mapper + : OMPRTL__tgt_target_data_end_mapper; break; case OMPD_target_update: - RTLFn = HasNowait ? OMPRTL__tgt_target_data_update_nowait - : OMPRTL__tgt_target_data_update; + RTLFn = HasNowait ? OMPRTL__tgt_target_data_update_nowait_mapper + : OMPRTL__tgt_target_data_update_mapper; break; case OMPD_parallel: case OMPD_for: @@ -10209,17 +10552,19 @@ MappableExprsHandler::MapValuesArrayTy Pointers; MappableExprsHandler::MapValuesArrayTy Sizes; MappableExprsHandler::MapFlagsArrayTy MapTypes; + MappableExprsHandler::MapMappersArrayTy Mappers; // Get map clause information. MappableExprsHandler MEHandler(D, CGF); - MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes); + MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes, Mappers); TargetDataInfo Info; // Fill up the arrays and create the arguments. - emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Info); + emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, Mappers, + Info); emitOffloadingArraysArgument(CGF, Info.BasePointersArray, Info.PointersArray, Info.SizesArray, - Info.MapTypesArray, Info); + Info.MapTypesArray, Info.MappersArray, Info); InputInfo.NumberOfTargetItems = Info.NumberOfPtrs; InputInfo.BasePointersArray = Address(Info.BasePointersArray, CGM.getPointerAlign()); @@ -10227,6 +10572,7 @@ Address(Info.PointersArray, CGM.getPointerAlign()); InputInfo.SizesArray = Address(Info.SizesArray, CGM.getPointerAlign()); + InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign()); MapTypesArray = Info.MapTypesArray; if (D.hasClausesOfKind()) CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo); Index: lib/CodeGen/CGStmtOpenMP.cpp =================================================================== --- lib/CodeGen/CGStmtOpenMP.cpp +++ lib/CodeGen/CGStmtOpenMP.cpp @@ -3137,36 +3137,41 @@ VarDecl *BPVD = nullptr; VarDecl *PVD = nullptr; VarDecl *SVD = nullptr; + VarDecl *MVD = nullptr; if (InputInfo.NumberOfTargetItems > 0) { auto *CD = CapturedDecl::Create( getContext(), getContext().getTranslationUnitDecl(), /*NumParams=*/0); llvm::APInt ArrSize(/*numBits=*/32, InputInfo.NumberOfTargetItems); - QualType BaseAndPointersType = getContext().getConstantArrayType( + QualType BaseAndPointerAndMapperType = getContext().getConstantArrayType( getContext().VoidPtrTy, ArrSize, ArrayType::Normal, /*IndexTypeQuals=*/0); BPVD = createImplicitFirstprivateForType( - getContext(), Data, BaseAndPointersType, CD, S.getBeginLoc()); + getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc()); PVD = createImplicitFirstprivateForType( - getContext(), Data, BaseAndPointersType, CD, S.getBeginLoc()); + getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc()); QualType SizesType = getContext().getConstantArrayType( getContext().getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/1), ArrSize, ArrayType::Normal, /*IndexTypeQuals=*/0); SVD = createImplicitFirstprivateForType(getContext(), Data, SizesType, CD, S.getBeginLoc()); + MVD = createImplicitFirstprivateForType( + getContext(), Data, BaseAndPointerAndMapperType, CD, S.getBeginLoc()); TargetScope.addPrivate( BPVD, [&InputInfo]() { return InputInfo.BasePointersArray; }); TargetScope.addPrivate(PVD, [&InputInfo]() { return InputInfo.PointersArray; }); TargetScope.addPrivate(SVD, [&InputInfo]() { return InputInfo.SizesArray; }); + TargetScope.addPrivate(MVD, + [&InputInfo]() { return InputInfo.MappersArray; }); } (void)TargetScope.Privatize(); // Build list of dependences. for (const auto *C : S.getClausesOfKind()) for (const Expr *IRef : C->varlists()) Data.Dependences.emplace_back(C->getDependencyKind(), IRef); - auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD, + auto &&CodeGen = [&Data, &S, CS, &BodyGen, BPVD, PVD, SVD, MVD, &InputInfo](CodeGenFunction &CGF, PrePostActionTy &Action) { // Set proper addresses for generated private copies. OMPPrivateScope Scope(CGF); @@ -3207,6 +3212,8 @@ CGF.GetAddrOfLocalVar(PVD), /*Index=*/0); InputInfo.SizesArray = CGF.Builder.CreateConstArrayGEP( CGF.GetAddrOfLocalVar(SVD), /*Index=*/0); + InputInfo.MappersArray = CGF.Builder.CreateConstArrayGEP( + CGF.GetAddrOfLocalVar(MVD), /*Index=*/0); } Action.Enter(CGF); Index: lib/CodeGen/CodeGenFunction.h =================================================================== --- lib/CodeGen/CodeGenFunction.h +++ lib/CodeGen/CodeGenFunction.h @@ -3114,12 +3114,15 @@ Address BasePointersArray = Address::invalid(); Address PointersArray = Address::invalid(); Address SizesArray = Address::invalid(); + Address MappersArray = Address::invalid(); unsigned NumberOfTargetItems = 0; explicit OMPTargetDataInfo() = default; OMPTargetDataInfo(Address BasePointersArray, Address PointersArray, - Address SizesArray, unsigned NumberOfTargetItems) + Address SizesArray, Address MappersArray, + unsigned NumberOfTargetItems) : BasePointersArray(BasePointersArray), PointersArray(PointersArray), - SizesArray(SizesArray), NumberOfTargetItems(NumberOfTargetItems) {} + SizesArray(SizesArray), MappersArray(MappersArray), + NumberOfTargetItems(NumberOfTargetItems) {} }; void EmitOMPTargetTaskBasedDirective(const OMPExecutableDirective &S, const RegionCodeGenTy &BodyGen, Index: test/OpenMP/capturing_in_templates.cpp =================================================================== --- test/OpenMP/capturing_in_templates.cpp +++ test/OpenMP/capturing_in_templates.cpp @@ -18,7 +18,7 @@ // CHECK-LABEL: @main int main(int argc, char **argv) { -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{.+}}.region_id, i32 0, i8** null, i8** null, i64* null, i64* null) +// CHECK: call i32 @__tgt_target_mapper(i64 -1, i8* @{{.+}}.region_id, i32 0, i8** null, i8** null, i64* null, i64* null, i8** null) #pragma omp target { for (int i = 0; i < 64; ++i) { Index: test/OpenMP/declare_mapper_codegen.cpp =================================================================== --- test/OpenMP/declare_mapper_codegen.cpp +++ test/OpenMP/declare_mapper_codegen.cpp @@ -22,14 +22,11 @@ #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: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 1] // 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: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 1] // 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: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 1] // CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34] class C { @@ -40,7 +37,7 @@ #pragma omp declare mapper(id: C s) map(s.a, s.b[0:2]) -// CK0-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C.id{{.*}}(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}) +// CK0: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}) // CK0: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]] // CK0: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]] // CK0: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]] @@ -209,41 +206,53 @@ 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: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** [[MPRGEP:%.+]]) // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 + // CK0-DAG: [[MPRGEP]] = getelementptr inbounds {{.+}}[[MPR:%[^,]+]], 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: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i32 0, i32 0 // CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C** // CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C** + // CK0-DAG: [[CMPR1:%.+]] = bitcast i8** [[MPR1]] to void (i8*, i8*, i8*, i64, i64)** // CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]] // CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]] + // CK0-DAG: store void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]], void (i8*, i8*, i8*, i64, i64)** [[CMPR1]] // CK0: call void [[KERNEL:@.+]](%class.C* [[VAL]]) #pragma omp target map(mapper(id),tofrom: c) { - ++c.a; + ++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: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}, i8** [[TMPRGEP:%.+]]) // CK0-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 // CK0-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[TMPRGEP]] = getelementptr inbounds {{.+}}[[TMPR:%[^,]+]], i32 0, i32 0 // CK0-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0 // CK0-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i32 0, i32 0 // CK0-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C** // CK0-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C** + // CK0-DAG: [[TCMPR1:%.+]] = bitcast i8** [[TMPR1]] to void (i8*, i8*, i8*, i64, i64)** // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]] // CK0-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]] + // CK0-DAG: store void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]], void (i8*, i8*, i8*, i64, i64)** [[TCMPR1]] #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: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}, i8** [[FMPRGEP:%.+]]) // CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 // CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[FMPRGEP]] = getelementptr inbounds {{.+}}[[FMPR:%[^,]+]], i32 0, i32 0 // CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0 // CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0 + // CK0-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i32 0, i32 0 // CK0-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C** // CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C** + // CK0-DAG: [[FCMPR1:%.+]] = bitcast i8** [[FMPR1]] to void (i8*, i8*, i8*, i64, i64)** // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]] // CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]] + // CK0-DAG: store void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]], void (i8*, i8*, i8*, i64, i64)** [[FCMPR1]] #pragma omp target update from(mapper(id): c) } @@ -257,7 +266,7 @@ // CK0: {{.+}} = add nsw i32 [[VAL]], 1 // CK0: } -#endif +#endif // CK0 ///==========================================================================/// @@ -276,6 +285,7 @@ // 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 +// C++ template template class C { @@ -409,6 +419,168 @@ // CK1: [[DONE]] // CK1: ret void -#endif +#endif // CK1 -#endif + +///==========================================================================/// +// RUN: %clang_cc1 -DCK2 -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 CK2 --check-prefix CK2-64 %s +// RUN: %clang_cc1 -DCK2 -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 -DCK2 -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 CK2 --check-prefix CK2-64 %s +// RUN: %clang_cc1 -DCK2 -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 CK2 --check-prefix CK2-32 %s +// RUN: %clang_cc1 -DCK2 -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 -DCK2 -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 CK2 --check-prefix CK2-32 %s + +// RUN: %clang_cc1 -DCK2 -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 -DCK2 -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 -DCK2 -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 -DCK2 -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 -DCK2 -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 -DCK2 -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 CK2 +// Nested mappers. + +class B { +public: + double a; +}; + +class C { +public: + double a; + B b; +}; + +#pragma omp declare mapper(B s) map(s.a) + +#pragma omp declare mapper(id: C s) map(s.b) + +// CK2: define {{.*}}void [[BMPRFUNC:@[.]omp_mapper[.].*B[.]default]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}) + +// CK2-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id(i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}) +// CK2: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]] +// CK2: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]] +// CK2: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]] +// CK2: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]] +// CK2: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]] +// CK2-DAG: [[SIZE:%.+]] = load i64, i64* [[SIZEADDR]] +// CK2-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] +// CK2-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]] +// CK2-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** +// CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]] +// CK2-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]] +// CK2-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]] +// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 +// CK2: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] + +// CK2: [[INITEVALDEL]] +// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK2: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK2: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] +// CK2: [[INIT]] +// CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 +// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]]) +// CK2: br label %[[LHEAD:[^,]+]] + +// CK2: [[LHEAD]] +// CK2: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK2: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]] +// CK2: [[LBODY]] +// CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK2: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] +// CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK2-DAG: [[BEND:%.+]] = getelementptr %class.B, %class.B* [[BBEGIN]], i32 1 +// CK2-DAG: [[BBEGINV:%.+]] = bitcast %class.B* [[BBEGIN]] to i8* +// CK2-DAG: [[BENDV:%.+]] = bitcast %class.B* [[BEND]] to i8* +// CK2-DAG: [[BBEGINI:%.+]] = ptrtoint i8* [[BBEGINV]] to i64 +// CK2-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64 +// CK2-DAG: [[BSIZE:%.+]] = sub i64 [[BENDI]], [[BBEGINI]] +// CK2-DAG: [[BUSIZE:%.+]] = sdiv exact i64 [[BSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK2-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* +// CK2-DAG: [[PTRADDR0BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8* +// CK2-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) +// CK2-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 +// CK2-DAG: br label %[[MEMBER:[^,]+]] +// CK2-DAG: [[MEMBER]] +// CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] +// CK2-DAG: [[MEMBERCOM]] +// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]] +// CK2-DAG: br label %[[LTYPE]] +// CK2-DAG: [[LTYPE]] +// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK2-DAG: [[ALLOC]] +// CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 +// CK2-DAG: br label %[[TYEND:[^,]+]] +// CK2-DAG: [[ALLOCELSE]] +// CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK2-DAG: [[TO]] +// CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 +// CK2-DAG: br label %[[TYEND]] +// CK2-DAG: [[TOELSE]] +// CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK2-DAG: [[FROM]] +// CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 +// CK2-DAG: br label %[[TYEND]] +// CK2-DAG: [[TYEND]] +// CK2-DAG: [[TYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] +// CK2-64: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[BUSIZE]], i64 [[TYPE0]]) +// CK2-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* +// CK2-DAG: [[PTRADDR1BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8* +// CK2-DAG: br label %[[MEMBER:[^,]+]] +// CK2-DAG: [[MEMBER]] +// CK2-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] +// CK2-DAG: [[MEMBERCOM]] +// 281474976710659 == 0x1,000,000,003 +// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] +// CK2-DAG: br label %[[LTYPE]] +// CK2-DAG: [[LTYPE]] +// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK2-DAG: [[ALLOC]] +// CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 +// CK2-DAG: br label %[[TYEND:[^,]+]] +// CK2-DAG: [[ALLOCELSE]] +// CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK2-DAG: [[TO]] +// CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 +// CK2-DAG: br label %[[TYEND]] +// CK2-DAG: [[TOELSE]] +// CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK2-DAG: [[FROM]] +// CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 +// CK2-DAG: br label %[[TYEND]] +// CK2-DAG: [[TYEND]] +// CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] +// CK2: call void [[BMPRFUNC]](i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 1, i64 [[TYPE1]]) +// CK2: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 +// CK2: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] + +// CK2: [[LEXIT]] +// CK2: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 +// CK2: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] +// CK2: [[EVALDEL]] +// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK2: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK2: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] +// CK2: [[DEL]] +// CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 +// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK2: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]]) +// CK2: br label %[[DONE]] +// CK2: [[DONE]] +// CK2: ret void + +#endif // CK2 + +#endif // HEADER Index: test/OpenMP/declare_target_link_codegen.cpp =================================================================== --- test/OpenMP/declare_target_link_codegen.cpp +++ test/OpenMP/declare_target_link_codegen.cpp @@ -77,9 +77,9 @@ // HOST: [[BP0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // HOST: [[P0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 -// HOST: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0)) +// HOST: call i32 @__tgt_target_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0)) // HOST: call void @__omp_offloading_{{.*}}_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* %{{[^,]+}}) -// HOST: call i32 @__tgt_target_teams(i64 -1, i8* @.__omp_offloading_{{.+}}_l47.region_id, i32 2, {{.+}}) +// HOST: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @.__omp_offloading_{{.+}}_l47.region_id, i32 2, {{.+}}) // HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* dereferenceable{{.*}}) // HOST: [[C:%.*]] = load i32, i32* @c, Index: test/OpenMP/target_is_device_ptr_codegen.cpp =================================================================== --- test/OpenMP/target_is_device_ptr_codegen.cpp +++ test/OpenMP/target_is_device_ptr_codegen.cpp @@ -49,7 +49,7 @@ float *l; T *t; - // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES00]]{{.+}}, {{.+}}[[TYPES00]]{{.+}}) + // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES00]]{{.+}}, {{.+}}[[TYPES00]]{{.+}}) // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 @@ -66,7 +66,7 @@ ++g; } - // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES01]]{{.+}}, {{.+}}[[TYPES01]]{{.+}}) + // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES01]]{{.+}}, {{.+}}[[TYPES01]]{{.+}}) // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 @@ -83,7 +83,7 @@ ++l; } - // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES02]]{{.+}}, {{.+}}[[TYPES02]]{{.+}}) + // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES02]]{{.+}}, {{.+}}[[TYPES02]]{{.+}}) // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 @@ -100,7 +100,7 @@ ++t; } - // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES03]]{{.+}}, {{.+}}[[TYPES03]]{{.+}}) + // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES03]]{{.+}}, {{.+}}[[TYPES03]]{{.+}}) // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 @@ -118,7 +118,7 @@ ++lr; } - // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES04]]{{.+}}, {{.+}}[[TYPES04]]{{.+}}) + // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES04]]{{.+}}, {{.+}}[[TYPES04]]{{.+}}) // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 @@ -136,7 +136,7 @@ ++tr; } - // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES05]]{{.+}}, {{.+}}[[TYPES05]]{{.+}}) + // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES05]]{{.+}}, {{.+}}[[TYPES05]]{{.+}}) // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 @@ -154,7 +154,7 @@ ++tr; } - // CK1-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES06]]{{.+}}, {{.+}}[[TYPES06]]{{.+}}) + // CK1-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES06]]{{.+}}, {{.+}}[[TYPES06]]{{.+}}) // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 @@ -231,7 +231,7 @@ void foo(double *&arg) { int *la = 0; - // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK2-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] @@ -246,7 +246,7 @@ a++; } - // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) + // CK2-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}) // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] @@ -261,7 +261,7 @@ b++; } - // CK2-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK2-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]