diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7429,6 +7429,7 @@ // S1 s; // double *p; // struct S2 *ps; + // int &ref; // } // S2 s; // S2 *ps; @@ -7472,6 +7473,14 @@ // 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 // @@ -7726,6 +7735,8 @@ .getCanonicalType() ->isAnyPointerType()) || I->getAssociatedExpression()->getType()->isAnyPointerType(); + bool IsMemberReference = EncounteredME && MapDecl && + MapDecl->getType()->isLValueReferenceType(); bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous; if (OASE) @@ -7744,13 +7755,30 @@ "Unexpected expression"); Address LB = Address::invalid(); + Address LowestElem = Address::invalid(); 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) { + Address Base = BP; + QualType BaseType = EncounteredME->getBase()->getType(); + if (EncounteredME->isArrow()) { + if (!IsExpressionFirstInfo || FirstPointerInComplexData) + Base = + CGF.EmitLoadOfPointer(Base, BaseType->castAs()); + BaseType = BaseType->getPointeeType(); + } + LValue BaseLVal = CGF.MakeAddrLValue(Base, BaseType); + LowestElem = CGF.EmitLValueForFieldInitialization( + BaseLVal, cast(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 @@ -7767,11 +7795,11 @@ "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 = { @@ -7797,10 +7825,28 @@ 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(VD); + if (FD && FD->getType()->isLValueReferenceType()) { + Address Base = BP; + const auto *ME = + cast(MC.getAssociatedExpression()); + QualType BaseType = ME->getBase()->getType(); + if (ME->isArrow()) { + if (!IsExpressionFirstInfo || FirstPointerInComplexData) + Base = CGF.EmitLoadOfPointer( + Base, BaseType->castAs()); + BaseType = BaseType->getPointeeType(); + } + LValue BaseLVal = CGF.MakeAddrLValue(Base, BaseType); + 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())); @@ -7856,10 +7902,10 @@ 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) @@ -7888,21 +7934,21 @@ // 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}; } } @@ -7916,7 +7962,7 @@ // The pointer becomes the base for the next element. if (Next != CE) - BP = LB; + LowestElem = BP = LB; IsExpressionFirstInfo = false; IsCaptureFirstInfo = false; diff --git a/clang/test/OpenMP/target_map_codegen_28.cpp b/clang/test/OpenMP/target_map_codegen_28.cpp --- a/clang/test/OpenMP/target_map_codegen_28.cpp +++ b/clang/test/OpenMP/target_map_codegen_28.cpp @@ -70,7 +70,7 @@ // 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]] @@ -116,11 +116,10 @@ // 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 @@ -128,9 +127,10 @@ // 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 [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CBP1]] // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]] // CK29-DAG: store i64 {{8|4}}, i64* [[S1]] + // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000]], // CK29-DAG: [[VAR1]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 0 // CK29-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 @@ -161,11 +161,10 @@ // 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 @@ -173,9 +172,10 @@ // 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 [[SSA]]** [[VAR00:%.+]], [[SSA]]*** [[CBP1]] // CK29-DAG: store double** [[VAR1:%.+]], double*** [[CP1]] // CK29-DAG: store i64 {{8|4}}, i64* [[S1]] + // CK29-DAG: [[VAR00]] = load [[SSA]]**, [[SSA]]*** [[VAR000]], // CK29-DAG: [[VAR1]] = load double**, double*** [[VAR1_REF:%.+]], // CK29-DAG: [[VAR1_REF]] = getelementptr inbounds [[SSA]], [[SSA]]* %{{.+}}, i32 0, i32 1 diff --git a/clang/test/OpenMP/target_map_codegen_35.cpp b/clang/test/OpenMP/target_map_codegen_35.cpp new file mode 100644 --- /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 --- /dev/null +++ b/openmp/libomptarget/test/mapping/data_member_ref.cpp @@ -0,0 +1,49 @@ +// 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 + +struct View { + int Data; +}; + +struct Foo { + Foo(View &V) : VRef(V) {} + View &VRef; +}; + +int main() { + View V; + V.Data = 123456; + Foo Bar(V); + + // 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); + return 0; +}