Skip to content

Commit

Permalink
[OpenMP] Codegen for target update directive.
Browse files Browse the repository at this point in the history 
Summary: This patch implements the code generation for the `target update` directive. The implemntation relies on the logic already in place for target data standalone directives, i.e. target enter/exit data.

Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev

Subscribers: caomhin, cfe-commits

Differential Revision: http://reviews.llvm.org/D20650

llvm-svn: 270886
  • Loading branch information
Samuel Antao committed May 26, 2016
1 parent 143f684 commit 8d2d730
Show file tree
Hide file tree
Showing 4 changed files with 331 additions and 27 deletions.
78 changes: 62 additions & 16 deletions clang/lib/CodeGen/CGOpenMPRuntime.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -660,6 +660,9 @@ enum OpenMPRTLFunction {
// Call to void __tgt_target_data_end(int32_t device_id, int32_t arg_num,
// void** args_base, void **args, size_t *arg_sizes, int32_t *arg_types);
OMPRTL__tgt_target_data_end,
// Call to void __tgt_target_data_update(int32_t device_id, int32_t arg_num,
// void** args_base, void **args, size_t *arg_sizes, int32_t *arg_types);
OMPRTL__tgt_target_data_update,
};

/// A basic class for pre|post-action for advanced codegen sequence for OpenMP
Expand Down Expand Up @@ -1609,6 +1612,20 @@ CGOpenMPRuntime::createRuntimeFunction(unsigned Function) {
RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_end");
break;
}
case OMPRTL__tgt_target_data_update: {
// Build void __tgt_target_data_update(int32_t device_id, int32_t arg_num,
// void** args_base, void **args, size_t *arg_sizes, int32_t *arg_types);
llvm::Type *TypeParams[] = {CGM.Int32Ty,
CGM.Int32Ty,
CGM.VoidPtrPtrTy,
CGM.VoidPtrPtrTy,
CGM.SizeTy->getPointerTo(),
CGM.Int32Ty->getPointerTo()};
llvm::FunctionType *FnTy =
llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update");
break;
}
}
assert(RTLFn && "Unable to find OpenMP runtime function");
return RTLFn;
Expand Down Expand Up @@ -5347,13 +5364,27 @@ class MappableExprsHandler {
// declaration in a single chunk so that we can generate the map flags
// correctly. Therefore, we organize all lists in a map.
llvm::DenseMap<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, OpenMPMapClauseKind MapModifier) {
const ValueDecl *VD =
D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
Info[VD].push_back({L, MapType, MapModifier});
};

for (auto *C : Directive.getClausesOfKind<OMPMapClause>())
for (auto L : C->component_lists()) {
const ValueDecl *VD =
L.first ? cast<ValueDecl>(L.first->getCanonicalDecl()) : nullptr;
Info[VD].push_back(
{L.second, C->getMapType(), C->getMapTypeModifier()});
}
for (auto L : C->component_lists())
InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifier());
for (auto *C : Directive.getClausesOfKind<OMPToClause>())
for (auto L : C->component_lists())
InfoGen(L.first, L.second, OMPC_MAP_to, OMPC_MAP_unknown);
for (auto *C : Directive.getClausesOfKind<OMPFromClause>())
for (auto L : C->component_lists())
InfoGen(L.first, L.second, OMPC_MAP_from, OMPC_MAP_unknown);

for (auto &M : Info) {
// We need to know when we generate information for the first component
Expand Down Expand Up @@ -6128,15 +6159,16 @@ void CGOpenMPRuntime::emitTargetDataCalls(CodeGenFunction &CGF,
}
}

void CGOpenMPRuntime::emitTargetEnterOrExitDataCall(
void CGOpenMPRuntime::emitTargetDataStandAloneCall(
CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *IfCond,
const Expr *Device) {
if (!CGF.HaveInsertPoint())
return;

assert((isa<OMPTargetEnterDataDirective>(D) ||
isa<OMPTargetExitDataDirective>(D)) &&
"Expecting either target enter or exit data directives.");
isa<OMPTargetExitDataDirective>(D) ||
isa<OMPTargetUpdateDirective>(D)) &&
"Expecting either target enter, exit data, or update directives.");

// Generate the code for the opening of the data environment.
auto &&ThenGen = [&D, &CGF, Device](CodeGenFunction &CGF, PrePostActionTy &) {
Expand All @@ -6147,8 +6179,8 @@ void CGOpenMPRuntime::emitTargetEnterOrExitDataCall(
MappableExprsHandler::MapFlagsArrayTy MapTypes;

// Get map clause information.
MappableExprsHandler MCHandler(D, CGF);
MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);
MappableExprsHandler MEHandler(D, CGF);
MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes);

llvm::Value *BasePointersArrayArg = nullptr;
llvm::Value *PointersArrayArg = nullptr;
Expand Down Expand Up @@ -6178,12 +6210,26 @@ void CGOpenMPRuntime::emitTargetEnterOrExitDataCall(
llvm::Value *OffloadingArgs[] = {
DeviceID, PointerNum, BasePointersArrayArg,
PointersArrayArg, SizesArrayArg, MapTypesArrayArg};

auto &RT = CGF.CGM.getOpenMPRuntime();
CGF.EmitRuntimeCall(
RT.createRuntimeFunction(isa<OMPTargetEnterDataDirective>(D)
? OMPRTL__tgt_target_data_begin
: OMPRTL__tgt_target_data_end),
OffloadingArgs);
// Select the right runtime function call for each expected standalone
// directive.
OpenMPRTLFunction RTLFn;
switch (D.getDirectiveKind()) {
default:
llvm_unreachable("Unexpected standalone target data directive.");
break;
case OMPD_target_enter_data:
RTLFn = OMPRTL__tgt_target_data_begin;
break;
case OMPD_target_exit_data:
RTLFn = OMPRTL__tgt_target_data_end;
break;
case OMPD_target_update:
RTLFn = OMPRTL__tgt_target_data_update;
break;
}
CGF.EmitRuntimeCall(RT.createRuntimeFunction(RTLFn), OffloadingArgs);
};

// In the event we get an if clause, we don't have to take any action on the
Expand Down
12 changes: 6 additions & 6 deletions clang/lib/CodeGen/CGOpenMPRuntime.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1009,17 +1009,17 @@ class CGOpenMPRuntime {
const Expr *IfCond, const Expr *Device,
const RegionCodeGenTy &CodeGen);

/// \brief Emit the target enter or exit data mapping code associated with
/// directive \a D.
/// \brief Emit the data mapping/movement code associated with the directive
/// \a D that should be of the form 'target [{enter|exit} data | update]'.
/// \param D Directive to emit.
/// \param IfCond Expression evaluated in if clause associated with the target
/// directive, or null if no if clause is used.
/// \param Device Expression evaluated in device clause associated with the
/// target directive, or null if no device clause is used.
virtual void emitTargetEnterOrExitDataCall(CodeGenFunction &CGF,
const OMPExecutableDirective &D,
const Expr *IfCond,
const Expr *Device);
virtual void emitTargetDataStandAloneCall(CodeGenFunction &CGF,
const OMPExecutableDirective &D,
const Expr *IfCond,
const Expr *Device);

/// Marks function \a Fn with properly mangled versions of vector functions.
/// \param FD Function marked as 'declare simd'.
Expand Down
23 changes: 18 additions & 5 deletions clang/lib/CodeGen/CGStmtOpenMP.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3366,8 +3366,7 @@ void CodeGenFunction::EmitOMPTargetEnterDataDirective(
if (auto *C = S.getSingleClause<OMPDeviceClause>())
Device = C->getDevice();

CGM.getOpenMPRuntime().emitTargetEnterOrExitDataCall(*this, S, IfCond,
Device);
CGM.getOpenMPRuntime().emitTargetDataStandAloneCall(*this, S, IfCond, Device);
}

void CodeGenFunction::EmitOMPTargetExitDataDirective(
Expand All @@ -3387,8 +3386,7 @@ void CodeGenFunction::EmitOMPTargetExitDataDirective(
if (auto *C = S.getSingleClause<OMPDeviceClause>())
Device = C->getDevice();

CGM.getOpenMPRuntime().emitTargetEnterOrExitDataCall(*this, S, IfCond,
Device);
CGM.getOpenMPRuntime().emitTargetDataStandAloneCall(*this, S, IfCond, Device);
}

void CodeGenFunction::EmitOMPTargetParallelDirective(
Expand Down Expand Up @@ -3550,5 +3548,20 @@ void CodeGenFunction::EmitOMPTaskLoopSimdDirective(
// Generate the instructions for '#pragma omp target update' directive.
void CodeGenFunction::EmitOMPTargetUpdateDirective(
const OMPTargetUpdateDirective &S) {
// TODO: codegen for target update
// If we don't have target devices, don't bother emitting the data mapping
// code.
if (CGM.getLangOpts().OMPTargetTriples.empty())
return;

// Check if we have any if clause associated with the directive.
const Expr *IfCond = nullptr;
if (auto *C = S.getSingleClause<OMPIfClause>())
IfCond = C->getCondition();

// Check if we have any device clause associated with the directive.
const Expr *Device = nullptr;
if (auto *C = S.getSingleClause<OMPDeviceClause>())
Device = C->getDevice();

CGM.getOpenMPRuntime().emitTargetDataStandAloneCall(*this, S, IfCond, Device);
}

0 comments on commit 8d2d730

Please sign in to comment.