[OpenMP 5.0] Codegen support for user-defined mappers.
This patch implements the code generation for OpenMP 5.0 declare mapper (user-defined mapper) constructs. For each declare mapper, a mapper function is generated. These mapper functions will be called by the runtime and/or other mapper functions to achieve user defined mapping. The design slides can be found at https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx Re-commit after revert in r367773 because r367755 changed the LLVM-IR output such that a CHECK line failed. Patch by Lingda Li <lildmh@gmail.com> Differential Revision: https://reviews.llvm.org/D59474 llvm-svn: 367905
This commit is contained in:
Michael Kruse
parent
37aa8ad663
commit
d47b9438d7
|
|
@ -59,6 +59,7 @@ 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)
|
||||
|
|
|
|||
|
|
@ -9860,7 +9860,7 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
|
|||
return !D->getDeclContext()->isDependentContext();
|
||||
else if (isa<OMPAllocateDecl>(D))
|
||||
return !D->getDeclContext()->isDependentContext();
|
||||
else if (isa<OMPDeclareReductionDecl>(D))
|
||||
else if (isa<OMPDeclareReductionDecl>(D) || isa<OMPDeclareMapperDecl>(D))
|
||||
return !D->getDeclContext()->isDependentContext();
|
||||
else if (isa<ImportDecl>(D))
|
||||
return true;
|
||||
|
|
|
|||
|
|
@ -2530,10 +2530,11 @@ void CodeGenModule::EmitOMPDeclareReduction(const OMPDeclareReductionDecl *D,
|
|||
}
|
||||
|
||||
void CodeGenModule::EmitOMPDeclareMapper(const OMPDeclareMapperDecl *D,
|
||||
CodeGenFunction *CGF) {
|
||||
if (!LangOpts.OpenMP || (!LangOpts.EmitAllDecls && !D->isUsed()))
|
||||
CodeGenFunction *CGF) {
|
||||
if (!LangOpts.OpenMP || LangOpts.OpenMPSimd ||
|
||||
(!LangOpts.EmitAllDecls && !D->isUsed()))
|
||||
return;
|
||||
// FIXME: need to implement mapper code generation
|
||||
getOpenMPRuntime().emitUserDefinedMapper(D, CGF);
|
||||
}
|
||||
|
||||
void CodeGenModule::EmitOMPRequiresDecl(const OMPRequiresDecl *D) {
|
||||
|
|
|
|||
|
|
@ -752,6 +752,11 @@ 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
|
||||
|
|
@ -1686,6 +1691,12 @@ 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() {
|
||||
|
|
@ -2459,6 +2470,24 @@ 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;
|
||||
|
|
@ -7093,6 +7122,15 @@ 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 {
|
||||
|
|
@ -7156,8 +7194,11 @@ private:
|
|||
: IE(IE), VD(VD) {}
|
||||
};
|
||||
|
||||
/// Directive from where the map clauses were extracted.
|
||||
const OMPExecutableDirective &CurDir;
|
||||
/// The target directive from where the mappable clauses were extracted. It
|
||||
/// is either a executable directive or a user-defined mapper directive.
|
||||
llvm::PointerUnion<const OMPExecutableDirective *,
|
||||
const OMPDeclareMapperDecl *>
|
||||
CurDir;
|
||||
|
||||
/// Function the directive is being generated for.
|
||||
CodeGenFunction &CGF;
|
||||
|
|
@ -7760,9 +7801,9 @@ private:
|
|||
}
|
||||
|
||||
static OpenMPOffloadMappingFlags getMemberOfFlag(unsigned Position) {
|
||||
// Member of is given by the 16 MSB of the flag, so rotate by 48 bits.
|
||||
// Rotate by getFlagMemberOffset() bits.
|
||||
return static_cast<OpenMPOffloadMappingFlags>(((uint64_t)Position + 1)
|
||||
<< 48);
|
||||
<< getFlagMemberOffset());
|
||||
}
|
||||
|
||||
static void setCorrectMemberOfFlag(OpenMPOffloadMappingFlags &Flags,
|
||||
|
|
@ -7842,7 +7883,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<OMPFirstprivateClause>())
|
||||
for (const auto *D : C->varlists())
|
||||
|
|
@ -7854,6 +7895,10 @@ 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.
|
||||
|
|
@ -7915,18 +7960,20 @@ public:
|
|||
IsImplicit);
|
||||
};
|
||||
|
||||
// FIXME: MSVC 2013 seems to require this-> to find member CurDir.
|
||||
for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>())
|
||||
assert(CurDir.is<const OMPExecutableDirective *>() &&
|
||||
"Expect a executable directive");
|
||||
const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
|
||||
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>())
|
||||
for (const auto &L : C->component_lists()) {
|
||||
InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifiers(),
|
||||
/*ReturnDevicePointer=*/false, C->isImplicit());
|
||||
}
|
||||
for (const auto *C : this->CurDir.getClausesOfKind<OMPToClause>())
|
||||
for (const auto *C : CurExecDir->getClausesOfKind<OMPToClause>())
|
||||
for (const auto &L : C->component_lists()) {
|
||||
InfoGen(L.first, L.second, OMPC_MAP_to, llvm::None,
|
||||
/*ReturnDevicePointer=*/false, C->isImplicit());
|
||||
}
|
||||
for (const auto *C : this->CurDir.getClausesOfKind<OMPFromClause>())
|
||||
for (const auto *C : CurExecDir->getClausesOfKind<OMPFromClause>())
|
||||
for (const auto &L : C->component_lists()) {
|
||||
InfoGen(L.first, L.second, OMPC_MAP_from, llvm::None,
|
||||
/*ReturnDevicePointer=*/false, C->isImplicit());
|
||||
|
|
@ -7941,9 +7988,8 @@ public:
|
|||
llvm::MapVector<const ValueDecl *, SmallVector<DeferredDevicePtrEntryTy, 4>>
|
||||
DeferredInfo;
|
||||
|
||||
// FIXME: MSVC 2013 seems to require this-> to find member CurDir.
|
||||
for (const auto *C :
|
||||
this->CurDir.getClausesOfKind<OMPUseDevicePtrClause>()) {
|
||||
CurExecDir->getClausesOfKind<OMPUseDevicePtrClause>()) {
|
||||
for (const auto &L : C->component_lists()) {
|
||||
assert(!L.second.empty() && "Not expecting empty list of components!");
|
||||
const ValueDecl *VD = L.second.back().getAssociatedDeclaration();
|
||||
|
|
@ -7972,7 +8018,6 @@ 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<MemberExpr>(IE)) {
|
||||
// Insert the pointer into Info to be processed by
|
||||
// generateInfoForComponentList. Because it is a member pointer
|
||||
|
|
@ -7985,11 +8030,11 @@ public:
|
|||
/*ReturnDevicePointer=*/false, C->isImplicit());
|
||||
DeferredInfo[nullptr].emplace_back(IE, VD);
|
||||
} else {
|
||||
llvm::Value *Ptr = this->CGF.EmitLoadOfScalar(
|
||||
this->CGF.EmitLValue(IE), IE->getExprLoc());
|
||||
llvm::Value *Ptr =
|
||||
CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
|
||||
BasePointers.emplace_back(Ptr, VD);
|
||||
Pointers.push_back(Ptr);
|
||||
Sizes.push_back(llvm::Constant::getNullValue(this->CGF.Int64Ty));
|
||||
Sizes.push_back(llvm::Constant::getNullValue(CGF.Int64Ty));
|
||||
Types.push_back(OMP_MAP_RETURN_PARAM | OMP_MAP_TARGET_PARAM);
|
||||
}
|
||||
}
|
||||
|
|
@ -8013,11 +8058,10 @@ public:
|
|||
|
||||
// Remember the current base pointer index.
|
||||
unsigned CurrentBasePointersIdx = CurBasePointers.size();
|
||||
// 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);
|
||||
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.
|
||||
|
|
@ -8069,6 +8113,78 @@ 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<const OMPDeclareMapperDecl *>() &&
|
||||
"Expect a declare mapper directive");
|
||||
const auto *CurMapperDir = CurDir.get<const OMPDeclareMapperDecl *>();
|
||||
// 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<const ValueDecl *, SmallVector<MapInfo, 8>> 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<OpenMPMapModifierKind> MapModifiers,
|
||||
bool ReturnDevicePointer, bool IsImplicit) {
|
||||
const ValueDecl *VD =
|
||||
D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
|
||||
Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
|
||||
IsImplicit);
|
||||
};
|
||||
|
||||
for (const auto *C : CurMapperDir->clauselists()) {
|
||||
const auto *MC = cast<OMPMapClause>(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,
|
||||
|
|
@ -8192,8 +8308,10 @@ public:
|
|||
std::tuple<OMPClauseMappableExprCommon::MappableExprComponentListRef,
|
||||
OpenMPMapClauseKind, ArrayRef<OpenMPMapModifierKind>, bool>;
|
||||
SmallVector<MapData, 4> DeclComponentLists;
|
||||
// FIXME: MSVC 2013 seems to require this-> to find member CurDir.
|
||||
for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>()) {
|
||||
assert(CurDir.is<const OMPExecutableDirective *>() &&
|
||||
"Expect a executable directive");
|
||||
const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
|
||||
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
|
||||
for (const auto &L : C->decl_component_lists(VD)) {
|
||||
assert(L.first == VD &&
|
||||
"We got information for the wrong declaration??");
|
||||
|
|
@ -8341,9 +8459,12 @@ public:
|
|||
MapValuesArrayTy &Pointers,
|
||||
MapValuesArrayTy &Sizes,
|
||||
MapFlagsArrayTy &Types) const {
|
||||
assert(CurDir.is<const OMPExecutableDirective *>() &&
|
||||
"Expect a executable directive");
|
||||
const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
|
||||
// Map other list items in the map clause which are not captured variables
|
||||
// but "declare target link" global variables.
|
||||
for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>()) {
|
||||
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
|
||||
for (const auto &L : C->component_lists()) {
|
||||
if (!L.first)
|
||||
continue;
|
||||
|
|
@ -8570,6 +8691,7 @@ 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(
|
||||
|
|
@ -8700,6 +8822,337 @@ 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.<type_name>.<mapper_id>.(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<VarDecl>(cast<DeclRefExpr>(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<PointerType>())
|
||||
.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<llvm::Value *(
|
||||
|
|
|
|||
|
|
@ -345,6 +345,14 @@ private:
|
|||
SmallVector<const OMPDeclareReductionDecl *, 4>>
|
||||
FunctionUDRMapTy;
|
||||
FunctionUDRMapTy FunctionUDRMap;
|
||||
/// Map from the user-defined mapper declaration to its corresponding
|
||||
/// functions.
|
||||
llvm::DenseMap<const OMPDeclareMapperDecl *, llvm::Function *> UDMMap;
|
||||
/// Map of functions and their local user-defined mappers.
|
||||
using FunctionUDMMapTy =
|
||||
llvm::DenseMap<llvm::Function *,
|
||||
SmallVector<const OMPDeclareMapperDecl *, 4>>;
|
||||
FunctionUDMMapTy FunctionUDMMap;
|
||||
/// Type kmp_critical_name, originally defined as typedef kmp_int32
|
||||
/// kmp_critical_name[8];
|
||||
llvm::ArrayType *KmpCriticalNameTy;
|
||||
|
|
@ -738,6 +746,14 @@ 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;
|
||||
|
|
@ -798,6 +814,10 @@ public:
|
|||
virtual std::pair<llvm::Function *, llvm::Function *>
|
||||
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*).
|
||||
|
|
|
|||
|
|
@ -232,6 +232,9 @@ namespace {
|
|||
if (auto *DRD = dyn_cast<OMPDeclareReductionDecl>(Member)) {
|
||||
if (Ctx->DeclMustBeEmitted(DRD))
|
||||
Builder->EmitGlobal(DRD);
|
||||
} else if (auto *DMD = dyn_cast<OMPDeclareMapperDecl>(Member)) {
|
||||
if (Ctx->DeclMustBeEmitted(DMD))
|
||||
Builder->EmitGlobal(DMD);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,92 +1,414 @@
|
|||
///==========================================================================///
|
||||
// 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)
|
||||
#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
|
||||
|
||||
// CHECK-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l54.region_id = weak constant i8 0
|
||||
// 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: [[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: [[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-LABEL: foo{{.*}}(
|
||||
// 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{{.*}}
|
||||
void foo(int a){
|
||||
int i = a;
|
||||
C c;
|
||||
c.a = a;
|
||||
|
||||
// CHECK-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
|
||||
// CHECK-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
|
||||
// CHECK-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
|
||||
// CHECK-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
|
||||
// CHECK-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
|
||||
// CHECK-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
|
||||
// CHECK-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
|
||||
// CHECK-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
|
||||
// CHECK-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
|
||||
// CHECK: call void [[KERNEL:@.+]](%class.C* [[VAL]])
|
||||
// CK0-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
|
||||
// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
|
||||
// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
|
||||
// CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
|
||||
// CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
|
||||
// CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
|
||||
// CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
|
||||
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
|
||||
// CK0: call void [[KERNEL:@.+]](%class.C* [[VAL]])
|
||||
#pragma omp target map(mapper(id),tofrom: c)
|
||||
{
|
||||
++c.a;
|
||||
}
|
||||
|
||||
// CHECK-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}})
|
||||
// CHECK-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
|
||||
// CHECK-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
|
||||
// CHECK-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CHECK-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CHECK-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C**
|
||||
// CHECK-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C**
|
||||
// CHECK-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]]
|
||||
// CHECK-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]]
|
||||
// CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], 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]]
|
||||
#pragma omp target update to(mapper(id): c)
|
||||
|
||||
// CHECK-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}})
|
||||
// CHECK-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
|
||||
// CHECK-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
|
||||
// CHECK-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CHECK-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CHECK-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
|
||||
// CHECK-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
|
||||
// CHECK-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
|
||||
// CHECK-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
|
||||
// CK0-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], 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]]
|
||||
#pragma omp target update from(mapper(id): c)
|
||||
}
|
||||
|
||||
|
||||
// CHECK: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]])
|
||||
// CHECK: [[ADDR:%.+]] = alloca %class.C*,
|
||||
// CHECK: store %class.C* [[ARG]], %class.C** [[ADDR]]
|
||||
// CHECK: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
|
||||
// CHECK: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
|
||||
// CHECK: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
|
||||
// CHECK: {{.+}} = add nsw i32 [[VAL]], 1
|
||||
// CHECK: }
|
||||
// CK0: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]])
|
||||
// CK0: [[ADDR:%.+]] = alloca %class.C*,
|
||||
// CK0: store %class.C* [[ARG]], %class.C** [[ADDR]]
|
||||
// CK0: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
|
||||
// CK0: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
|
||||
// CK0: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
|
||||
// CK0: {{.+}} = add nsw i32 [[VAL]], 1
|
||||
// CK0: }
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
|
||||
// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
|
||||
// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
|
||||
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
|
||||
// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
|
||||
// RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
|
||||
|
||||
// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
|
||||
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
|
||||
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
|
||||
// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
|
||||
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
|
||||
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
|
||||
|
||||
#ifdef CK1
|
||||
|
||||
template <class T>
|
||||
class C {
|
||||
public:
|
||||
T a;
|
||||
};
|
||||
|
||||
#pragma omp declare mapper(id: C<int> 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
|
||||
|
||||
#endif
|
||||
|
|
|
|||
Loading…
Reference in New Issue