[Openmp-commits] [openmp] 0411b23 - [OPENMP]Map data field with l-value reference types.
Alexey Bataev via Openmp-commits
openmp-commits at lists.llvm.org
Mon Mar 29 07:25:10 PDT 2021
Author: Alexey Bataev
Date: 2021-03-29T07:07:09-07:00
New Revision: 0411b2331916cc8c7a8be9dd0eb540b731e6d9ce
URL: https://github.com/llvm/llvm-project/commit/0411b2331916cc8c7a8be9dd0eb540b731e6d9ce
DIFF: https://github.com/llvm/llvm-project/commit/0411b2331916cc8c7a8be9dd0eb540b731e6d9ce.diff
LOG: [OPENMP]Map data field with l-value reference types.
Added initial support dfor the mapping of the data members with l-value
reference types.
Differential Revision: https://reviews.llvm.org/D98812
Added:
clang/test/OpenMP/target_map_codegen_35.cpp
openmp/libomptarget/test/mapping/data_member_ref.cpp
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
clang/test/OpenMP/target_map_codegen_28.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 466ff096b585..95859e6e94a7 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7439,6 +7439,7 @@ class MappableExprsHandler {
// S1 s;
// double *p;
// struct S2 *ps;
+ // int &ref;
// }
// S2 s;
// S2 *ps;
@@ -7482,6 +7483,14 @@ class MappableExprsHandler {
// optimizes this entry out, same in the examples below)
// (***) map the pointee (map: to)
//
+ // map(to: s.ref)
+ // &s, &(s.ref), sizeof(int*), TARGET_PARAM (*)
+ // &s, &(s.ref), sizeof(int), MEMBER_OF(1) | PTR_AND_OBJ | TO (***)
+ // (*) alloc space for struct members, only this is a target parameter
+ // (**) map the pointer (nothing to be mapped in this example) (the compiler
+ // optimizes this entry out, same in the examples below)
+ // (***) map the pointee (map: to)
+ //
// map(s.ps)
// &s, &(s.ps), sizeof(S2*), TARGET_PARAM | TO | FROM
//
@@ -7679,6 +7688,7 @@ class MappableExprsHandler {
uint64_t DimSize = 1;
bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous;
+ bool IsPrevMemberReference = false;
for (; I != CE; ++I) {
// If the current component is member of a struct (parent struct) mark it.
@@ -7736,12 +7746,16 @@ class MappableExprsHandler {
.getCanonicalType()
->isAnyPointerType()) ||
I->getAssociatedExpression()->getType()->isAnyPointerType();
+ bool IsMemberReference = isa<MemberExpr>(I->getAssociatedExpression()) &&
+ MapDecl &&
+ MapDecl->getType()->isLValueReferenceType();
bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous;
if (OASE)
++DimSize;
- if (Next == CE || IsNonDerefPointer || IsFinalArraySection) {
+ if (Next == CE || IsMemberReference || IsNonDerefPointer ||
+ IsFinalArraySection) {
// If this is not the last component, we expect the pointer to be
// associated with an array expression or member expression.
assert((Next == CE ||
@@ -7754,22 +7768,53 @@ class MappableExprsHandler {
"Unexpected expression");
Address LB = Address::invalid();
+ Address LowestElem = Address::invalid();
+ auto &&EmitMemberExprBase = [](CodeGenFunction &CGF,
+ const MemberExpr *E) {
+ const Expr *BaseExpr = E->getBase();
+ // If this is s.x, emit s as an lvalue. If it is s->x, emit s as a
+ // scalar.
+ LValue BaseLV;
+ if (E->isArrow()) {
+ LValueBaseInfo BaseInfo;
+ TBAAAccessInfo TBAAInfo;
+ Address Addr =
+ CGF.EmitPointerWithAlignment(BaseExpr, &BaseInfo, &TBAAInfo);
+ QualType PtrTy = BaseExpr->getType()->getPointeeType();
+ BaseLV = CGF.MakeAddrLValue(Addr, PtrTy, BaseInfo, TBAAInfo);
+ } else {
+ BaseLV = CGF.EmitOMPSharedLValue(BaseExpr);
+ }
+ return BaseLV;
+ };
if (OAShE) {
- LB = Address(CGF.EmitScalarExpr(OAShE->getBase()),
- CGF.getContext().getTypeAlignInChars(
- OAShE->getBase()->getType()));
- } else {
- LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
+ LowestElem = LB = Address(CGF.EmitScalarExpr(OAShE->getBase()),
+ CGF.getContext().getTypeAlignInChars(
+ OAShE->getBase()->getType()));
+ } else if (IsMemberReference) {
+ const auto *ME = cast<MemberExpr>(I->getAssociatedExpression());
+ LValue BaseLVal = EmitMemberExprBase(CGF, ME);
+ LowestElem = CGF.EmitLValueForFieldInitialization(
+ BaseLVal, cast<FieldDecl>(MapDecl))
+ .getAddress(CGF);
+ LB = CGF.EmitLoadOfReferenceLValue(LowestElem, MapDecl->getType())
.getAddress(CGF);
+ } else {
+ LowestElem = LB =
+ CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
+ .getAddress(CGF);
}
// If this component is a pointer inside the base struct then we don't
// need to create any entry for it - it will be combined with the object
// it is pointing to into a single PTR_AND_OBJ entry.
bool IsMemberPointerOrAddr =
- (IsPointer || ForDeviceAddr) && EncounteredME &&
- (dyn_cast<MemberExpr>(I->getAssociatedExpression()) ==
- EncounteredME);
+ EncounteredME &&
+ (((IsPointer || ForDeviceAddr) &&
+ I->getAssociatedExpression() == EncounteredME) ||
+ (IsPrevMemberReference && !IsPointer) ||
+ (IsMemberReference && Next != CE &&
+ !Next->getAssociatedExpression()->getType()->isPointerType()));
if (!OverlappedElements.empty() && Next == CE) {
// Handle base element with the info for overlapped elements.
assert(!PartialStruct.Base.isValid() && "The base element is set.");
@@ -7777,11 +7822,11 @@ class MappableExprsHandler {
"Unexpected base element with the pointer type.");
// Mark the whole struct as the struct that requires allocation on the
// device.
- PartialStruct.LowestElem = {0, LB};
+ PartialStruct.LowestElem = {0, LowestElem};
CharUnits TypeSize = CGF.getContext().getTypeSizeInChars(
I->getAssociatedExpression()->getType());
Address HB = CGF.Builder.CreateConstGEP(
- CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(LB,
+ CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(LowestElem,
CGF.VoidPtrTy),
TypeSize.getQuantity() - 1);
PartialStruct.HighestElem = {
@@ -7807,10 +7852,20 @@ class MappableExprsHandler {
Address ComponentLB = Address::invalid();
for (const OMPClauseMappableExprCommon::MappableComponent &MC :
Component) {
- if (MC.getAssociatedDeclaration()) {
- ComponentLB =
- CGF.EmitOMPSharedLValue(MC.getAssociatedExpression())
- .getAddress(CGF);
+ if (const ValueDecl *VD = MC.getAssociatedDeclaration()) {
+ const auto *FD = dyn_cast<FieldDecl>(VD);
+ if (FD && FD->getType()->isLValueReferenceType()) {
+ const auto *ME =
+ cast<MemberExpr>(MC.getAssociatedExpression());
+ LValue BaseLVal = EmitMemberExprBase(CGF, ME);
+ ComponentLB =
+ CGF.EmitLValueForFieldInitialization(BaseLVal, FD)
+ .getAddress(CGF);
+ } else {
+ ComponentLB =
+ CGF.EmitOMPSharedLValue(MC.getAssociatedExpression())
+ .getAddress(CGF);
+ }
Size = CGF.Builder.CreatePtrDiff(
CGF.EmitCastToVoidPtr(ComponentLB.getPointer()),
CGF.EmitCastToVoidPtr(LB.getPointer()));
@@ -7833,8 +7888,7 @@ class MappableExprsHandler {
CombinedInfo.BasePointers.push_back(BP.getPointer());
CombinedInfo.Pointers.push_back(LB.getPointer());
Size = CGF.Builder.CreatePtrDiff(
- CGF.EmitCastToVoidPtr(
- CGF.Builder.CreateConstGEP(HB, 1).getPointer()),
+ CGF.Builder.CreateConstGEP(HB, 1).getPointer(),
CGF.EmitCastToVoidPtr(LB.getPointer()));
CombinedInfo.Sizes.push_back(
CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
@@ -7866,13 +7920,13 @@ class MappableExprsHandler {
OpenMPOffloadMappingFlags Flags = getMapTypeBits(
MapType, MapModifiers, MotionModifiers, IsImplicit,
!IsExpressionFirstInfo || RequiresReference ||
- FirstPointerInComplexData,
+ FirstPointerInComplexData || IsMemberReference,
IsCaptureFirstInfo && !RequiresReference, IsNonContiguous);
- if (!IsExpressionFirstInfo) {
+ if (!IsExpressionFirstInfo || IsMemberReference) {
// If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
// then we reset the TO/FROM/ALWAYS/DELETE/CLOSE flags.
- if (IsPointer)
+ if (IsPointer || (IsMemberReference && Next != CE))
Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS |
OMP_MAP_DELETE | OMP_MAP_CLOSE);
@@ -7898,21 +7952,21 @@ class MappableExprsHandler {
// Update info about the lowest and highest elements for this struct
if (!PartialStruct.Base.isValid()) {
- PartialStruct.LowestElem = {FieldIndex, LB};
+ PartialStruct.LowestElem = {FieldIndex, LowestElem};
if (IsFinalArraySection) {
Address HB =
CGF.EmitOMPArraySectionExpr(OASE, /*IsLowerBound=*/false)
.getAddress(CGF);
PartialStruct.HighestElem = {FieldIndex, HB};
} else {
- PartialStruct.HighestElem = {FieldIndex, LB};
+ PartialStruct.HighestElem = {FieldIndex, LowestElem};
}
PartialStruct.Base = BP;
PartialStruct.LB = BP;
} else if (FieldIndex < PartialStruct.LowestElem.first) {
- PartialStruct.LowestElem = {FieldIndex, LB};
+ PartialStruct.LowestElem = {FieldIndex, LowestElem};
} else if (FieldIndex > PartialStruct.HighestElem.first) {
- PartialStruct.HighestElem = {FieldIndex, LB};
+ PartialStruct.HighestElem = {FieldIndex, LowestElem};
}
}
@@ -7926,11 +7980,12 @@ class MappableExprsHandler {
// The pointer becomes the base for the next element.
if (Next != CE)
- BP = LB;
+ BP = IsMemberReference ? LowestElem : LB;
IsExpressionFirstInfo = false;
IsCaptureFirstInfo = false;
FirstPointerInComplexData = false;
+ IsPrevMemberReference = IsMemberReference;
} else if (FirstPointerInComplexData) {
QualType Ty = Components.rbegin()
->getAssociatedDeclaration()
@@ -8914,7 +8969,6 @@ class MappableExprsHandler {
std::prev(It)
->getAssociatedExpression()
->getType()
- .getNonReferenceType()
->isPointerType())
continue;
const MapData &BaseData = CI == CE ? L : L1;
diff --git a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
index 73ef18651b52..dfeb350877e6 100644
--- a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
@@ -383,10 +383,10 @@ struct ST {
a++;
// CK2: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
- // CK2: [[CBP:%.+]] = bitcast i8** [[BP]] to double***
- // CK2: store double** [[RVAL:%.+]], double*** [[CBP]],
+ // CK2: [[CBP:%.+]] = bitcast i8** [[BP]] to double****
+ // CK2: store double*** [[RVAL:%.+]], double**** [[CBP]],
// CK2: call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
- // CK2: [[CBP1:%.+]] = bitcast double*** [[CBP]] to double**
+ // CK2: [[CBP1:%.+]] = bitcast double**** [[CBP]] to double**
// CK2: [[VAL:%.+]] = load double*, double** [[CBP1]],
// CK2: store double* [[VAL]], double** [[PVT:%.+]],
// CK2: store double** [[PVT]], double*** [[PVT2:%.+]],
@@ -428,8 +428,8 @@ struct ST {
la++;
// CK2: [[BP1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 1
- // CK2: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double***
- // CK2: store double** [[RVAL1:%.+]], double*** [[CBP1]],
+ // CK2: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double****
+ // CK2: store double*** [[RVAL1:%.+]], double**** [[CBP1]],
// CK2: [[BP2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2
// CK2: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double***
// CK2: store double** [[RVAL2:%.+]], double*** [[CBP2]],
@@ -438,7 +438,7 @@ struct ST {
// CK2: [[VAL2:%.+]] = load double*, double** [[_CBP2]],
// CK2: store double* [[VAL2]], double** [[PVT2:%.+]],
// CK2: store double** [[PVT2]], double*** [[_PVT2:%.+]],
- // CK2: [[_CBP1:%.+]] = bitcast double*** [[CBP1]] to double**
+ // CK2: [[_CBP1:%.+]] = bitcast double**** [[CBP1]] to double**
// CK2: [[VAL1:%.+]] = load double*, double** [[_CBP1]],
// CK2: store double* [[VAL1]], double** [[PVT1:%.+]],
// CK2: store double** [[PVT1]], double*** [[_PVT1:%.+]],
diff --git a/clang/test/OpenMP/target_map_codegen_28.cpp b/clang/test/OpenMP/target_map_codegen_28.cpp
index 32fcfc9e439f..9d561eb5a0e5 100644
--- a/clang/test/OpenMP/target_map_codegen_28.cpp
+++ b/clang/test/OpenMP/target_map_codegen_28.cpp
@@ -37,13 +37,13 @@
// CK29: [[SSB:%.+]] = type { [[SSA]]*, [[SSA]]** }
// CK29-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
-// CK29: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710672, i64 19]
+// CK29: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675]
// CK29-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK29: [[MTYPE01:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710672, i64 19]
// CK29-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
-// CK29: [[MTYPE02:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710672, i64 19]
+// CK29: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675]
struct SSA{
double *p;
@@ -60,7 +60,7 @@ struct SSB{
void foo() {
// Region 00
- // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null, i8** null)
+ // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[Z:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null, i8** null)
// CK29-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK29-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -70,32 +70,22 @@ struct SSB{
// CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
- // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**
+ // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]***
// CK29-DAG: store [[SSB]]* [[VAR0:%.+]], [[SSB]]** [[CBP0]]
// CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]]
// CK29-DAG: store i64 %{{.+}}, i64* [[S0]]
// CK29-DAG: [[VAR0]] = load [[SSB]]*, [[SSB]]** %
// CK29-DAG: [[VAR00]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 0
- // CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK29-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
- // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]***
- // CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
- // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]]
- // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]]
- // CK29-DAG: store i64 {{8|4}}, i64* [[S1]]
- // CK29-DAG: [[VAR1]] = load double**, double*** [[VAR1_REF:%.+]],
- // CK29-DAG: [[VAR1_REF]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1
-
- // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
- // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
- // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
- // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double***
+ // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+ // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double****
// CK29-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double**
- // CK29-DAG: store double** [[VAR1]], double*** [[CBP2]]
+ // CK29-DAG: store double*** [[VAR1:%.+]], double**** [[CBP2]]
// CK29-DAG: store double* [[VAR2:%.+]], double** [[CP2]]
// CK29-DAG: store i64 80, i64* [[S2]]
+ // CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1
// CK29-DAG: [[VAR2]] = getelementptr inbounds double, double* [[VAR22:%.+]], i{{.+}} 0
// CK29-DAG: [[VAR22]] = load double*, double** %{{.+}},
@@ -116,19 +106,18 @@ struct SSB{
// CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
- // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**
+ // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]****
// CK29-DAG: store [[SSB]]* [[VAR0]], [[SSB]]** [[CBP0]]
- // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]]
+ // CK29-DAG: store [[SSA]]*** [[VAR000:%.+]], [[SSA]]**** [[CP0]]
// CK29-DAG: store i64 %{{.+}}, i64* [[S0]]
- // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000:%.+]],
// CK29-DAG: [[VAR000]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 1
// CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
// CK29-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
// CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
- // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]***
+ // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]****
// CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
- // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]]
+ // CK29-DAG: store [[SSA]]*** [[VAR000]], [[SSA]]**** [[CBP1]]
// CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]]
// CK29-DAG: store i64 {{8|4}}, i64* [[S1]]
// CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 0
@@ -151,7 +140,7 @@ struct SSB{
}
// Region 02
- // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null, i8** null)
+ // CK29-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null, i8** null)
// CK29-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK29-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -161,32 +150,21 @@ struct SSB{
// CK29-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK29-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
- // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]**
+ // CK29-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SSA]]****
// CK29-DAG: store [[SSB]]* [[VAR0]], [[SSB]]** [[CBP0]]
- // CK29-DAG: store [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CP0]]
+ // CK29-DAG: store [[SSA]]*** [[VAR000:%.+]], [[SSA]]**** [[CP0]]
// CK29-DAG: store i64 %{{.+}}, i64* [[S0]]
- // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000:%.+]],
// CK29-DAG: [[VAR000]] = getelementptr inbounds [[SSB]], [[SSB]]* [[VAR0]], i32 0, i32 1
- // CK29-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK29-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK29-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
- // CK29-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]***
- // CK29-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double***
- // CK29-DAG: store [[SSA]]** [[VAR00]], [[SSA]]*** [[CBP1]]
- // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]]
- // CK29-DAG: store i64 {{8|4}}, i64* [[S1]]
- // CK29-DAG: [[VAR1]] = load double**, double*** [[VAR1_REF:%.+]],
- // CK29-DAG: [[VAR1_REF]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1
-
- // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
- // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
- // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
- // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double***
+ // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK29-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK29-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+ // CK29-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double****
// CK29-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double**
- // CK29-DAG: store double** [[VAR1]], double*** [[CBP2]]
+ // CK29-DAG: store double*** [[VAR1:%.+]], double**** [[CBP2]]
// CK29-DAG: store double* [[VAR2:%.+]], double** [[CP2]]
// CK29-DAG: store i64 80, i64* [[S2]]
+ // CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1
// CK29-DAG: [[VAR2]] = getelementptr inbounds double, double* [[VAR22:%.+]], i{{.+}} 0
// CK29-DAG: [[VAR22]] = load double*, double** %{{.+}},
diff --git a/clang/test/OpenMP/target_map_codegen_35.cpp b/clang/test/OpenMP/target_map_codegen_35.cpp
new file mode 100644
index 000000000000..74fbcf703958
--- /dev/null
+++ b/clang/test/OpenMP/target_map_codegen_35.cpp
@@ -0,0 +1,182 @@
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK35 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-64
+// RUN: %clang_cc1 -DCK35 -fopenmp -fopenmp-version=50 -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-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-64
+// RUN: %clang_cc1 -DCK35 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-32
+// RUN: %clang_cc1 -DCK35 -fopenmp -fopenmp-version=50 -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-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK35 --check-prefix CK35-32
+
+// RUN: %clang_cc1 -DCK35 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK35 -fopenmp-simd -fopenmp-version=50 -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-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK35 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// RUN: %clang_cc1 -DCK35 -fopenmp-simd -fopenmp-version=50 -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-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
+// SIMD-ONLY32-NOT: {{__kmpc|__tgt}}
+#ifdef CK35
+
+class S {
+public:
+ S(double &b) : b(b) {}
+ int a;
+ double &b;
+ void foo();
+};
+
+// TARGET_PARAM = 0x20
+// MEMBER_OF_1 | TO = 0x1000000000001
+// MEMBER_OF_1 | PTR_AND_OBJ | TO = 0x1000000000011
+// CK35-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000011]]]
+// TARGET_PARAM = 0x20
+// MEMBER_OF_1 | PTR_AND_OBJ | FROM = 0x1000000000012
+// CK35-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [2 x i64] [i64 [[#0x20]], i64 [[#0x1000000000012]]]
+
+void ref_map() {
+ double b;
+ S s(b);
+
+ // CK35-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null, i8** null)
+ // CK35-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK35-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK35-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // pass TARGET_PARAM {&s, &s, ((void*)(&s+1)-(void*)&s)}
+
+ // CK35-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK35-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK35-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+ // CK35-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S**
+ // CK35-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to %class.S**
+
+ // CK35-DAG: store %class.S* [[S_ADDR:%.+]], %class.S** [[BPC0]],
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC0]],
+ // CK35-DAG: store i64 [[S_SIZE:%.+]], i64* [[S0]],
+
+ // CK35-DAG: [[S_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK35-DAG: [[SZ]] = sub i64 [[S_1_INTPTR:%.+]], [[S_INTPTR:%.+]]
+ // CK35-DAG: [[S_1_INTPTR]] = ptrtoint i8* [[S_1_VOID:%.+]] to i64
+ // CK35-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+ // CK35-DAG: [[S_1_VOID]] = bitcast %class.S* [[S_1:%.+]] to i8*
+ // CK35-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+ // CK35-DAG: [[S_1]] = getelementptr %class.S, %class.S* [[S_ADDR]], i32 1
+
+ // pass MEMBER_OF_1 | TO {&s, &s, ((void*)(&s.a+1)-(void*)&s)} to copy the data of s.a.
+
+ // CK35-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK35-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK35-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+
+ // CK35-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S**
+ // CK35-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to %class.S**
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]],
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[PC1]],
+ // CK35-DAG: store i64 [[A_SIZE:%.+]], i64* [[S1]],
+
+ // CK35-DAG: [[A_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK35-DAG: [[SZ]] = sub i64 [[B_BEGIN_INTPTR:%.+]], [[S_INTPTR:%.+]]
+ // CK35-DAG: [[S_INTPTR]] = ptrtoint i8* [[S_VOID:%.+]] to i64
+ // CK35-DAG: [[B_BEGIN_INTPTR]] = ptrtoint i8* [[B_BEGIN_VOID:%.+]] to i64
+ // CK35-DAG: [[S_VOID]] = bitcast %class.S* [[S_ADDR]] to i8*
+ // CK35-DAG: [[B_BEGIN_VOID]] = bitcast double** [[B_ADDR:%.+]] to i8*
+ // CK35-DAG: [[B_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+ // pass MEMBER_OF_1 | TO {&s, &s.b+1, ((void*)(&s+1)-(void*)(&s.b+1))} to copy the data of remainder of s.
+
+ // CK35-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+ // CK35-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+ // CK35-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+
+ // CK35-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to %class.S**
+ // CK35-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double***
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC2]],
+ // CK35-DAG: store double** [[B_END:%.+]], double*** [[PC2]],
+ // CK35-DAG: store i64 [[REM_SIZE:%.+]], i64* [[S2]],
+
+ // CK35-DAG: [[B_END]] = getelementptr double*, double** [[B_ADDR]], i{{.+}} 1
+
+ // CK35-DAG: [[REM_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK35-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[B_END_INTPTR:%.+]]
+ // CK35-DAG: [[B_END_INTPTR]] = ptrtoint i8* [[B_END_VOID:%.+]] to i64
+ // CK35-DAG: [[S_END_INTPTR]] = ptrtoint i8* [[S_END_VOID:%.+]] to i64
+ // CK35-DAG: [[B_END_VOID]] = bitcast double** [[B_END]] to i8*
+ // CK35-DAG: [[S_END_VOID]] = getelementptr i8, i8* [[S_LAST:%.+]], i{{.+}} 1
+ // CK35-64-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOIDPTR:%.+]], i64 15
+ // CK35-32-DAG: [[S_LAST]] = getelementptr i8, i8* [[S_VOIDPTR:%.+]], i32 7
+ // CK35-DAG: [[S_VOIDPTR]] = bitcast %class.S* [[S_ADDR]] to i8*
+
+ // pass MEMBER_OF_1 | PTR_AND_OBJ | TO {&s, &s.b, 8|4} to copy the data of s.b.
+
+ // CK35-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+ // CK35-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+ // CK35-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+
+ // CK35-DAG: [[BPC3:%.+]] = bitcast i8** [[BP3]] to %class.S**
+ // CK35-DAG: [[PC3:%.+]] = bitcast i8** [[P3]] to double**
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC3]],
+ // CK35-DAG: store double* [[B_ADDR:%.+]], double** [[PC3]],
+ // CK35-DAG: store i64 8, i64* [[S3]],
+
+ // CK35-DAG: [[B_ADDR]] = load double*, double** [[B_REF:%.+]],
+ // CK35-DAG: [[B_REF]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+
+ #pragma omp target map(to: s, s.b)
+ s.foo();
+
+ // CK35 : call void
+
+ // CK35-DAG: call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null, i8** null)
+ // CK35-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK35-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+ // CK35-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+ // pass TARGET_PARAM {&s, &s.b, ((void*)(&s.b+1)-(void*)&s.b)}
+
+ // CK35-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK35-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK35-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+
+ // CK35-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to %class.S**
+ // CK35-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double***
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC0]],
+ // CK35-DAG: store double** [[SB_ADDR:%.+]], double*** [[PC0]],
+ // CK35-DAG: store i64 [[B_SIZE:%.+]], i64* [[S0]],
+
+ // CK35-DAG: [[B_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+ // CK35-DAG: [[SZ]] = sub i64 [[SB_1_INTPTR:%.+]], [[SB_INTPTR:%.+]]
+ // CK35-DAG: [[SB_1_INTPTR]] = ptrtoint i8* [[SB_1_VOID:%.+]] to i64
+ // CK35-DAG: [[SB_INTPTR]] = ptrtoint i8* [[SB_VOID:%.+]] to i64
+ // CK35-DAG: [[SB_1_VOID]] = bitcast double** [[SB_1:%.+]] to i8*
+ // CK35-DAG: [[SB_VOID]] = bitcast double** [[SB_ADDR:%.+]] to i8*
+ // CK35-DAG: [[SB_ADDR]] = getelementptr inbounds %class.S, %class.S* [[S_ADDR]], i32 0, i32 1
+ // CK35-DAG: [[SB_1]] = getelementptr double*, double** [[SB_ADDR]], i{{.+}} 1
+
+ // pass MEMBER_OF_1 | PTR_AND_OBJ | FROM {&s, &s.b, 8|4} to copy the data of s.c.
+
+ // CK35-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK35-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK35-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+
+ // CK35-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to %class.S**
+ // CK35-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double**
+
+ // CK35-DAG: store %class.S* [[S_ADDR]], %class.S** [[BPC1]],
+ // CK35-DAG: store double* [[B_ADDR:%.+]], double** [[PC1]],
+ // CK35-DAG: store i64 8, i64* [[S1]],
+
+ // CK35-DAG: [[B_ADDR]] = load double*, double** [[SB_ADDR]],
+
+ #pragma omp target map(from: s.b)
+ s.foo();
+}
+
+#endif // CK35
+#endif
diff --git a/openmp/libomptarget/test/mapping/data_member_ref.cpp b/openmp/libomptarget/test/mapping/data_member_ref.cpp
new file mode 100644
index 000000000000..44e4d1e71708
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/data_member_ref.cpp
@@ -0,0 +1,69 @@
+// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
+
+#include <stdio.h>
+
+struct View {
+ int Data;
+};
+
+struct ViewPtr {
+ int *Data;
+};
+
+template <typename T> struct Foo {
+ Foo(T &V) : VRef(V) {}
+ T &VRef;
+};
+
+int main() {
+ View V;
+ V.Data = 123456;
+ Foo<View> Bar(V);
+ ViewPtr V1;
+ int Data = 123456;
+ V1.Data = &Data;
+ Foo<ViewPtr> Baz(V1);
+
+ // CHECK: Host 123456.
+ printf("Host %d.\n", Bar.VRef.Data);
+#pragma omp target map(Bar.VRef)
+ {
+ // CHECK: Device 123456.
+ printf("Device %d.\n", Bar.VRef.Data);
+ V.Data = 654321;
+ // CHECK: Device 654321.
+ printf("Device %d.\n", Bar.VRef.Data);
+ }
+ // CHECK: Host 654321 654321.
+ printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
+ V.Data = 123456;
+ // CHECK: Host 123456.
+ printf("Host %d.\n", Bar.VRef.Data);
+#pragma omp target map(Bar) map(Bar.VRef)
+ {
+ // CHECK: Device 123456.
+ printf("Device %d.\n", Bar.VRef.Data);
+ V.Data = 654321;
+ // CHECK: Device 654321.
+ printf("Device %d.\n", Bar.VRef.Data);
+ }
+ // CHECK: Host 654321 654321.
+ printf("Host %d %d.\n", Bar.VRef.Data, V.Data);
+ // CHECK: Host 123456.
+ printf("Host %d.\n", *Baz.VRef.Data);
+#pragma omp target map(*Baz.VRef.Data)
+ {
+ // CHECK: Device 123456.
+ printf("Device %d.\n", *Baz.VRef.Data);
+ *V1.Data = 654321;
+ // CHECK: Device 654321.
+ printf("Device %d.\n", *Baz.VRef.Data);
+ }
+ // CHECK: Host 654321 654321 654321.
+ printf("Host %d %d %d.\n", *Baz.VRef.Data, *V1.Data, Data);
+ return 0;
+}
More information about the Openmp-commits
mailing list