/// 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,
};
/// \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();
}
}
+ /// \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.
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 {
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();
} 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!");
// 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;
{
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
projects / clang / commitdiff