Skip to content

Commit

Permalink
[OpenMP] Add support for the 'private pointer' flag to signal variabl…
Browse files Browse the repository at this point in the history 
…es captured in target regions and used in first-private clauses.

Summary: If a variable is implicitly mapped (doesn't show in a map clause), the runtime library has to be informed if the corresponding capture shows up in first-private clause, so that the storage previously allocated in the device is used. This patch adds the support for that.

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

Subscribers: caomhin, cfe-commits

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

llvm-svn: 270870
  • Loading branch information
Samuel Antao committed May 26, 2016
1 parent 6c42e06 commit d486f84
Show file tree
Hide file tree
Showing 3 changed files with 220 additions and 78 deletions.
188 changes: 117 additions & 71 deletions clang/lib/CodeGen/CGOpenMPRuntime.cpp
View file
Expand Up @@ -4927,6 +4927,9 @@ class MappableExprsHandler {
/// map/privatization results in multiple arguments passed to the runtime
/// library.
OMP_MAP_FIRST_REF = 0x20,
/// \brief This flag signals that the reference being passed is a pointer to
/// private data.
OMP_MAP_PRIVATE_PTR = 0x80,
/// \brief Pass the element to the device by value.
OMP_MAP_PRIVATE_VAL = 0x100,
};
Expand All @@ -4941,6 +4944,9 @@ class MappableExprsHandler {
/// \brief Function the directive is being generated for.
CodeGenFunction &CGF;

/// \brief Set of all first private variables in the current directive.
llvm::SmallPtrSet<const VarDecl *, 8> FirstPrivateDecls;

llvm::Value *getExprTypeSize(const Expr *E) const {
auto ExprTy = E->getType().getCanonicalType();

Expand Down Expand Up @@ -5293,9 +5299,33 @@ class MappableExprsHandler {
}
}

/// \brief Return the adjusted map modifiers if the declaration a capture
/// refers to appears in a first-private clause. This is expected to be used
/// only with directives that start with 'target'.
unsigned adjustMapModifiersForPrivateClauses(const CapturedStmt::Capture &Cap,
unsigned CurrentModifiers) {
assert(Cap.capturesVariable() && "Expected capture by reference only!");

// A first private variable captured by reference will use only the
// 'private ptr' and 'map to' flag. Return the right flags if the captured
// declaration is known as first-private in this handler.
if (FirstPrivateDecls.count(Cap.getCapturedVar()))
return MappableExprsHandler::OMP_MAP_PRIVATE_PTR |
MappableExprsHandler::OMP_MAP_TO;

// We didn't modify anything.
return CurrentModifiers;
}

public:
MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
: Directive(Dir), CGF(CGF) {}
: Directive(Dir), CGF(CGF) {
// Extract firstprivate clause information.
for (const auto *C : Dir.getClausesOfKind<OMPFirstprivateClause>())
for (const auto *D : C->varlists())
FirstPrivateDecls.insert(
cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl())->getCanonicalDecl());
}

/// \brief Generate all the base pointers, section pointers, sizes and map
/// types for the extracted mappable expressions.
Expand Down Expand Up @@ -5377,6 +5407,86 @@ class MappableExprsHandler {

return;
}

/// \brief Generate the default map information for a given capture \a CI,
/// record field declaration \a RI and captured value \a CV.
void generateDefaultMapInfo(
const CapturedStmt::Capture &CI, const FieldDecl &RI, llvm::Value *CV,
MappableExprsHandler::MapValuesArrayTy &CurBasePointers,
MappableExprsHandler::MapValuesArrayTy &CurPointers,
MappableExprsHandler::MapValuesArrayTy &CurSizes,
MappableExprsHandler::MapFlagsArrayTy &CurMapTypes) {
auto &Ctx = CGF.getContext();

// Do the default mapping.
if (CI.capturesThis()) {
CurBasePointers.push_back(CV);
CurPointers.push_back(CV);
const PointerType *PtrTy = cast<PointerType>(RI.getType().getTypePtr());
CurSizes.push_back(CGF.getTypeSize(PtrTy->getPointeeType()));
// Default map type.
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_TO |
MappableExprsHandler::OMP_MAP_FROM);
} else if (CI.capturesVariableByCopy()) {
if (!RI.getType()->isAnyPointerType()) {
// If the field is not a pointer, we need to save the actual value
// and load it as a void pointer.
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL);
auto DstAddr = CGF.CreateMemTemp(Ctx.getUIntPtrType(),
Twine(CI.getCapturedVar()->getName()) +
".casted");
LValue DstLV = CGF.MakeAddrLValue(DstAddr, Ctx.getUIntPtrType());

auto *SrcAddrVal = CGF.EmitScalarConversion(
DstAddr.getPointer(), Ctx.getPointerType(Ctx.getUIntPtrType()),
Ctx.getPointerType(RI.getType()), SourceLocation());
LValue SrcLV = CGF.MakeNaturalAlignAddrLValue(SrcAddrVal, RI.getType());

// Store the value using the source type pointer.
CGF.EmitStoreThroughLValue(RValue::get(CV), SrcLV);

// Load the value using the destination type pointer.
CurBasePointers.push_back(
CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal());
CurPointers.push_back(CurBasePointers.back());

// Get the size of the type to be used in the map.
CurSizes.push_back(CGF.getTypeSize(RI.getType()));
} else {
// Pointers are implicitly mapped with a zero size and no flags
// (other than first map that is added for all implicit maps).
CurMapTypes.push_back(0u);
CurBasePointers.push_back(CV);
CurPointers.push_back(CV);
CurSizes.push_back(llvm::Constant::getNullValue(CGF.SizeTy));
}
} else {
assert(CI.capturesVariable() && "Expected captured reference.");
CurBasePointers.push_back(CV);
CurPointers.push_back(CV);

const ReferenceType *PtrTy =
cast<ReferenceType>(RI.getType().getTypePtr());
QualType ElementType = PtrTy->getPointeeType();
CurSizes.push_back(CGF.getTypeSize(ElementType));
// The default map type for a scalar/complex type is 'to' because by
// default the value doesn't have to be retrieved. For an aggregate
// type, the default is 'tofrom'.
CurMapTypes.push_back(ElementType->isAggregateType()
? (MappableExprsHandler::OMP_MAP_TO |
MappableExprsHandler::OMP_MAP_FROM)
: MappableExprsHandler::OMP_MAP_TO);

// If we have a capture by reference we may need to add the private
// pointer flag if the base declaration shows in some first-private
// clause.
CurMapTypes.back() =
adjustMapModifiersForPrivateClauses(CI, CurMapTypes.back());
}
// Every default map produces a single argument, so, it is always the
// first one.
CurMapTypes.back() |= MappableExprsHandler::OMP_MAP_FIRST_REF;
}
};

enum OpenMPOffloadingReservedDeviceIDs {
Expand Down Expand Up @@ -5559,8 +5669,8 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
MappableExprsHandler::MapValuesArrayTy CurSizes;
MappableExprsHandler::MapFlagsArrayTy CurMapTypes;

// Get map clause information.
MappableExprsHandler MCHandler(D, CGF);
// Get mappable expression information.
MappableExprsHandler MEHandler(D, CGF);

const CapturedStmt &CS = *cast<CapturedStmt>(D.getAssociatedStmt());
auto RI = CS.getCapturedRecordDecl()->field_begin();
Expand Down Expand Up @@ -5588,75 +5698,11 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
} else {
// If we have any information in the map clause, we use it, otherwise we
// just do a default mapping.
MCHandler.generateInfoForCapture(CI, CurBasePointers, CurPointers,
MEHandler.generateInfoForCapture(CI, CurBasePointers, CurPointers,
CurSizes, CurMapTypes);

if (CurBasePointers.empty()) {
// Do the default mapping.
if (CI->capturesThis()) {
CurBasePointers.push_back(*CV);
CurPointers.push_back(*CV);
const PointerType *PtrTy =
cast<PointerType>(RI->getType().getTypePtr());
CurSizes.push_back(CGF.getTypeSize(PtrTy->getPointeeType()));
// Default map type.
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_TO |
MappableExprsHandler::OMP_MAP_FROM);
} else if (CI->capturesVariableByCopy()) {
if (!RI->getType()->isAnyPointerType()) {
// If the field is not a pointer, we need to save the actual value
// and load it as a void pointer.
CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL);
auto DstAddr = CGF.CreateMemTemp(
Ctx.getUIntPtrType(),
Twine(CI->getCapturedVar()->getName()) + ".casted");
LValue DstLV = CGF.MakeAddrLValue(DstAddr, Ctx.getUIntPtrType());

auto *SrcAddrVal = CGF.EmitScalarConversion(
DstAddr.getPointer(), Ctx.getPointerType(Ctx.getUIntPtrType()),
Ctx.getPointerType(RI->getType()), SourceLocation());
LValue SrcLV =
CGF.MakeNaturalAlignAddrLValue(SrcAddrVal, RI->getType());

// Store the value using the source type pointer.
CGF.EmitStoreThroughLValue(RValue::get(*CV), SrcLV);

// Load the value using the destination type pointer.
CurBasePointers.push_back(
CGF.EmitLoadOfLValue(DstLV, SourceLocation()).getScalarVal());
CurPointers.push_back(CurBasePointers.back());

// Get the size of the type to be used in the map.
CurSizes.push_back(CGF.getTypeSize(RI->getType()));
} else {
// Pointers are implicitly mapped with a zero size and no flags
// (other than first map that is added for all implicit maps).
CurMapTypes.push_back(0u);
CurBasePointers.push_back(*CV);
CurPointers.push_back(*CV);
CurSizes.push_back(llvm::Constant::getNullValue(CGM.SizeTy));
}
} else {
assert(CI->capturesVariable() && "Expected captured reference.");
CurBasePointers.push_back(*CV);
CurPointers.push_back(*CV);

const ReferenceType *PtrTy =
cast<ReferenceType>(RI->getType().getTypePtr());
QualType ElementType = PtrTy->getPointeeType();
CurSizes.push_back(CGF.getTypeSize(ElementType));
// The default map type for a scalar/complex type is 'to' because by
// default the value doesn't have to be retrieved. For an aggregate
// type, the default is 'tofrom'.
CurMapTypes.push_back(ElementType->isAggregateType()
? (MappableExprsHandler::OMP_MAP_TO |
MappableExprsHandler::OMP_MAP_FROM)
: MappableExprsHandler::OMP_MAP_TO);
}
// Every default map produces a single argument, so, it is always the
// first one.
CurMapTypes.back() |= MappableExprsHandler::OMP_MAP_FIRST_REF;
}
if (CurBasePointers.empty())
MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurBasePointers,
CurPointers, CurSizes, CurMapTypes);
}
// We expect to have at least an element of information for this capture.
assert(!CurBasePointers.empty() && "Non-existing map pointer for capture!");
Expand Down
8 changes: 4 additions & 4 deletions clang/test/OpenMP/target_firstprivate_codegen.cpp
View file
Expand Up @@ -34,14 +34,14 @@ struct TT{

// CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ:32|64]] 4]
// CHECK: [[MAPT:@.+]] = private unnamed_addr constant [1 x i32] [i32 288]
// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 35, i32 288, i32 35, i32 35, i32 288, i32 288, i32 35, i32 35]
// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 161, i32 288, i32 161, i32 161, i32 288, i32 288, i32 161, i32 161]
// CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] zeroinitializer
// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [1 x i32] [i32 32]
// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 35]
// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 161]
// CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 1, i[[SZ]] 40]
// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 35]
// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 161]
// CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 40]
// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 35]
// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 161]


// CHECK: define {{.*}}[[FOO:@.+]](
Expand Down
102 changes: 99 additions & 3 deletions clang/test/OpenMP/target_map_codegen.cpp
View file
Expand Up @@ -4281,8 +4281,17 @@ int explicit_maps_with_private_class_members(){
// CK27: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
// CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35]

// CK27-LABEL: zero_size_section_maps
void zero_size_section_maps (int ii){
// CK27: [[SIZE05:@.+]] = private {{.*}}constant [1 x i[[Z]]] zeroinitializer
// CK27: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i32] [i32 32]

// CK27: [[SIZE07:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4]
// CK27: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i32] [i32 288]

// CK27: [[SIZE09:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 40]
// CK27: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i32] [i32 161]

// CK27-LABEL: zero_size_section_and_private_maps
void zero_size_section_and_private_maps (int ii){

// Map of a pointer.
int *pa;
Expand Down Expand Up @@ -4367,12 +4376,99 @@ void zero_size_section_maps (int ii){
{
pa[50]++;
}

int *pvtPtr;
int pvtScl;
int pvtArr[10];

// Region 04
// CK27: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i32* null)
// CK27: call void [[CALL04:@.+]]()
#pragma omp target private(pvtPtr)
{
pvtPtr[5]++;
}

// Region 05
// CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}})
// CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
// CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
// CK27-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
// CK27-DAG: [[CPVAL0]] = bitcast i32* [[VAR0]] to i8*

// CK27: call void [[CALL05:@.+]](i32* {{[^,]+}})
#pragma omp target firstprivate(pvtPtr)
{
pvtPtr[5]++;
}

// Region 06
// CK27: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i32* null)
// CK27: call void [[CALL06:@.+]]()
#pragma omp target private(pvtScl)
{
pvtScl++;
}

// Region 07
// CK27-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZE07]]{{.+}}, {{.+}}[[MTYPE07]]{{.+}})
// CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK27-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
// CK27-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
// CK27-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
// CK27-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
// CK27-DAG: [[VALBP]] = inttoptr i[[Z]] [[VAL:%.+]] to i8*
// CK27-DAG: [[VALP]] = inttoptr i[[Z]] [[VAL:%.+]] to i8*
// CK27-DAG: [[VAL]] = load i[[Z]], i[[Z]]* [[ADDR:%.+]],
// CK27-64-DAG: [[CADDR:%.+]] = bitcast i[[Z]]* [[ADDR]] to i32*
// CK27-64-DAG: store i32 {{.+}}, i32* [[CADDR]],

// CK27: call void [[CALL07:@.+]](i[[Z]] [[VAL]])
#pragma omp target firstprivate(pvtScl)
{
pvtScl++;
}

// Region 08
// CK27: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 0, i8** null, i8** null, i{{64|32}}* null, i32* null)
// CK27: call void [[CALL08:@.+]]()
#pragma omp target private(pvtArr)
{
pvtArr[5]++;
}

// Region 09
// CK27-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE09]]{{.+}})
// CK27-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK27-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

// CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK27-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
// CK27-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
// CK27-DAG: [[CBPVAL0]] = bitcast [10 x i32]* [[VAR0:%.+]] to i8*
// CK27-DAG: [[CPVAL0]] = bitcast [10 x i32]* [[VAR0]] to i8*

// CK27: call void [[CALL09:@.+]]([10 x i32]* {{[^,]+}})
#pragma omp target firstprivate(pvtArr)
{
pvtArr[5]++;
}
}

// CK27: define {{.+}}[[CALL00]]
// CK27: define {{.+}}[[CALL01]]
// CK27: define {{.+}}[[CALL02]]
// CK27: define {{.+}}[[CALL03]]

// CK27: define {{.+}}[[CALL04]]
// CK27: define {{.+}}[[CALL05]]
// CK27: define {{.+}}[[CALL06]]
// CK27: define {{.+}}[[CALL07]]
#endif
#endif

0 comments on commit d486f84

Please sign in to comment.