maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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
This commit is contained in:
Alexey Bataev 2021-03-09 05:16:02 -08:00
parent d0d92fee6f
commit 0411b23319
5 changed files with 360 additions and 77 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

106
clang/lib/CodeGen/CGOpenMPRuntime.cpp
View File

@ -7439,6 +7439,7 @@ private:
// S1 s;
// double *p;
// struct S2 *ps;
// int &ref;
// }
// S2 s;
// S2 *ps;
@ -7482,6 +7483,14 @@ private:
// 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 @@ private:
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 @@ private:
.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 @@ private:
"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 @@ private:
"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 @@ private:
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 @@ private:
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 @@ private:
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 @@ private:
// 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 @@ private:
// 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 @@ public:
std::prev(It)
->getAssociatedExpression()
->getType()
.getNonReferenceType()
->isPointerType())
continue;
const MapData &BaseData = CI == CE ? L : L1;

12
clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
View File

@ -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:%.+]],

68
clang/test/OpenMP/target_map_codegen_28.cpp
View File

@ -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** %{{.+}},

182
clang/test/OpenMP/target_map_codegen_35.cpp Normal file
View File

@ -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

69
openmp/libomptarget/test/mapping/data_member_ref.cpp Normal file
View File

@ -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;
}