[Openmp-commits] [openmp] 622e461 - [OPENMP]Fix PR46824: Global declare target pointer cannot be accessed in target region.
Alexey Bataev via Openmp-commits
openmp-commits at lists.llvm.org
Thu Jul 30 08:19:36 PDT 2020
Author: Alexey Bataev
Date: 2020-07-30T11:18:33-04:00
New Revision: 622e46156d9a91206c877a604d069bb3e2dbf294
URL: https://github.com/llvm/llvm-project/commit/622e46156d9a91206c877a604d069bb3e2dbf294
DIFF: https://github.com/llvm/llvm-project/commit/622e46156d9a91206c877a604d069bb3e2dbf294.diff
LOG: [OPENMP]Fix PR46824: Global declare target pointer cannot be accessed in target region.
Need to map the base pointer for all directives, not only target
data-based ones.
The base pointer is mapped for array sections, array subscript, array
shaping and other array-like constructs with the base pointer. Also,
codegen for use_device_ptr clause was modified to correctly handle
mapping combination of array like constructs + use_device_ptr clause.
The data for use_device_ptr clause is emitted as the last records in the
data mapping array.
Reviewed By: ye-luo
Differential Revision: https://reviews.llvm.org/D84767
Added:
openmp/libomptarget/test/env/base_ptr_ref_count.c
Modified:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/target_data_codegen.cpp
clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
clang/test/OpenMP/target_map_codegen.cpp
clang/test/OpenMP/target_update_codegen.cpp
openmp/libomptarget/src/omptarget.cpp
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index b8e33948c21c..dc12286c72be 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7392,10 +7392,9 @@ class MappableExprsHandler {
// &p, &p, sizeof(float*), TARGET_PARAM | TO | FROM
//
// map(p[1:24])
+ // &p, &p[1], 24*sizeof(float), TARGET_PARAM | TO | FROM | PTR_AND_OBJ
+ // in unified shared memory mode or for local pointers
// p, &p[1], 24*sizeof(float), TARGET_PARAM | TO | FROM
- // for data directives
- // p, p, sizeof(float*), TARGET_PARAM | TO | FROM
- // p, &p[1], 24*sizeof(float), PTR_AND_OBJ | TO | FROM
//
// map(s)
// &s, &s, sizeof(S2), TARGET_PARAM | TO | FROM
@@ -7530,6 +7529,7 @@ class MappableExprsHandler {
// Track if the map information being generated is the first for a list of
// components.
bool IsExpressionFirstInfo = true;
+ bool FirstPointerInComplexData = false;
Address BP = Address::invalid();
const Expr *AssocExpr = I->getAssociatedExpression();
const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
@@ -7572,17 +7572,16 @@ class MappableExprsHandler {
QualType Ty =
I->getAssociatedDeclaration()->getType().getNonReferenceType();
if (Ty->isAnyPointerType() && std::next(I) != CE) {
- BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
-
- // For non-data directives, we do not need to generate individual map
- // information for the pointer, it can be associated with the combined
- // storage.
+ // No need to generate individual map information for the pointer, it
+ // can be associated with the combined storage if shared memory mode is
+ // active or the base declaration is not global variable.
+ const auto *VD = dyn_cast<VarDecl>(I->getAssociatedDeclaration());
if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
- !CurDir.is<const OMPExecutableDirective *>() ||
- !isOpenMPTargetDataManagementDirective(
- CurDir.get<const OMPExecutableDirective *>()
- ->getDirectiveKind()))
- ++I;
+ !VD || VD->hasLocalStorage())
+ BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
+ else
+ FirstPointerInComplexData = true;
+ ++I;
}
}
@@ -7617,8 +7616,19 @@ class MappableExprsHandler {
EncounteredME = dyn_cast<MemberExpr>(I->getAssociatedExpression());
// If we encounter a PTR_AND_OBJ entry from now on it should be marked
// as MEMBER_OF the parent struct.
- if (EncounteredME)
+ if (EncounteredME) {
ShouldBeMemberOf = true;
+ // Do not emit as complex pointer if this is actually not array-like
+ // expression.
+ if (FirstPointerInComplexData) {
+ QualType Ty = std::prev(I)
+ ->getAssociatedDeclaration()
+ ->getType()
+ .getNonReferenceType();
+ BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
+ FirstPointerInComplexData = false;
+ }
+ }
}
auto Next = std::next(I);
@@ -7760,7 +7770,8 @@ class MappableExprsHandler {
// (there is a set of entries for each capture).
OpenMPOffloadMappingFlags Flags =
getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit,
- !IsExpressionFirstInfo || RequiresReference,
+ !IsExpressionFirstInfo || RequiresReference ||
+ FirstPointerInComplexData,
IsCaptureFirstInfo && !RequiresReference);
if (!IsExpressionFirstInfo) {
@@ -7819,6 +7830,7 @@ class MappableExprsHandler {
IsExpressionFirstInfo = false;
IsCaptureFirstInfo = false;
+ FirstPointerInComplexData = false;
}
}
}
@@ -8067,6 +8079,7 @@ class MappableExprsHandler {
// emission of that entry until the whole struct has been processed.
llvm::MapVector<const ValueDecl *, SmallVector<DeferredDevicePtrEntryTy, 4>>
DeferredInfo;
+ MapCombinedInfoTy UseDevicePtrCombinedInfo;
for (const auto *C :
CurExecDir->getClausesOfKind<OMPUseDevicePtrClause>()) {
@@ -8086,15 +8099,27 @@ class MappableExprsHandler {
// We potentially have map information for this declaration already.
// Look for the first set of components that refer to it.
if (It != Info.end()) {
- auto CI = std::find_if(
- It->second.begin(), It->second.end(), [VD](const MapInfo &MI) {
- return MI.Components.back().getAssociatedDeclaration() == VD;
- });
+ auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) {
+ return MI.Components.back().getAssociatedDeclaration() == VD;
+ });
// If we found a map entry, signal that the pointer has to be returned
// and move on to the next declaration.
+ // Exclude cases where the base pointer is mapped as array subscript,
+ // array section or array shaping. The base address is passed as a
+ // pointer to base in this case and cannot be used as a base for
+ // use_device_ptr list item.
if (CI != It->second.end()) {
- CI->ReturnDevicePointer = true;
- continue;
+ auto PrevCI = std::next(CI->Components.rbegin());
+ const auto *VarD = dyn_cast<VarDecl>(VD);
+ if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
+ isa<MemberExpr>(IE) ||
+ !VD->getType().getNonReferenceType()->isPointerType() ||
+ PrevCI == CI->Components.rend() ||
+ isa<MemberExpr>(PrevCI->getAssociatedExpression()) || !VarD ||
+ VarD->hasLocalStorage()) {
+ CI->ReturnDevicePointer = true;
+ continue;
+ }
}
}
@@ -8115,13 +8140,13 @@ class MappableExprsHandler {
} else {
llvm::Value *Ptr =
CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc());
- CombinedInfo.BasePointers.emplace_back(Ptr, VD);
- CombinedInfo.Pointers.push_back(Ptr);
- CombinedInfo.Sizes.push_back(
+ UseDevicePtrCombinedInfo.BasePointers.emplace_back(Ptr, VD);
+ UseDevicePtrCombinedInfo.Pointers.push_back(Ptr);
+ UseDevicePtrCombinedInfo.Sizes.push_back(
llvm::Constant::getNullValue(CGF.Int64Ty));
- CombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM |
- OMP_MAP_TARGET_PARAM);
- CombinedInfo.Mappers.push_back(nullptr);
+ UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM |
+ OMP_MAP_TARGET_PARAM);
+ UseDevicePtrCombinedInfo.Mappers.push_back(nullptr);
}
}
}
@@ -8273,6 +8298,8 @@ class MappableExprsHandler {
// We need to append the results of this capture to what we already have.
CombinedInfo.append(CurInfo);
}
+ // Append data for use_device_ptr clauses.
+ CombinedInfo.append(UseDevicePtrCombinedInfo);
}
/// Generate all the base pointers, section pointers, sizes, map types, and
diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp
index 80e674c0b49d..a2bb8cdf5ba8 100644
--- a/clang/test/OpenMP/target_data_codegen.cpp
+++ b/clang/test/OpenMP/target_data_codegen.cpp
@@ -555,7 +555,7 @@ struct S2 {
void test_close_modifier(int arg) {
S2 *ps;
- // CK5: private unnamed_addr constant [6 x i64] [i64 1059, i64 32, i64 562949953422339, i64 562949953421328, i64 16, i64 1043]
+ // CK5: private unnamed_addr constant [5 x i64] [i64 1059, i64 32, i64 562949953421328, i64 16, i64 1043]
#pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s)
{
++(arg);
@@ -634,20 +634,17 @@ void test_present_modifier(int arg) {
// Make sure the struct picks up present even if another element of the struct
// doesn't have present.
- // CK8: private unnamed_addr constant [15 x i64]
+ // CK8: private unnamed_addr constant [11 x i64]
// ps1
//
// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
- // MEMBER_OF_1=0x1000000000000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1000000000013
- // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
//
// CK8-SAME: {{^}} [i64 [[#0x1020]], i64 [[#0x1000000000003]],
- // CK8-SAME: {{^}} i64 [[#0x1000000000013]], i64 [[#0x1000000001003]],
// CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]],
// arg
@@ -659,16 +656,13 @@ void test_present_modifier(int arg) {
// ps2
//
// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
- // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x9000000001003
- // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x9000000001013
- // MEMBER_OF_9=0x9000000000000 | FROM=0x2 | TO=0x1 = 0x9000000000003
- // MEMBER_OF_9=0x9000000000000 | PTR_AND_OBJ=0x10 = 0x9000000000010
+ // MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
+ // MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
// PTR_AND_OBJ=0x10 = 0x10
// PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
//
- // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x9000000001003]],
- // CK8-SAME: {{^}} i64 [[#0x9000000001013]], i64 [[#0x9000000000003]],
- // CK8-SAME: {{^}} i64 [[#0x9000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
+ // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x7000000001003]],
+ // CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
#pragma omp target data map(tofrom: ps1->s) \
map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \
map(tofrom: ps2->ps->ps->ps->s)
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 fe6ea01b43c9..ca5536f927a1 100644
--- a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
@@ -22,18 +22,18 @@
double *g;
// CK1: @g = global double*
-// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE03:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE04:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE05:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE06:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE07:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19]
-// CK1: [[MTYPE08:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 35, i64 19]
-// CK1: [[MTYPE09:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 99, i64 19]
-// CK1: [[MTYPE10:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 99, i64 19]
-// CK1: [[MTYPE11:@.+]] = {{.*}}constant [3 x i64] [i64 96, i64 35, i64 19]
-// CK1: [[MTYPE12:@.+]] = {{.*}}constant [3 x i64] [i64 96, i64 35, i64 19]
+// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
+// CK1: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 99]
+// CK1: [[MTYPE03:@.+]] = {{.*}}constant [1 x i64] [i64 99]
+// CK1: [[MTYPE04:@.+]] = {{.*}}constant [1 x i64] [i64 99]
+// CK1: [[MTYPE05:@.+]] = {{.*}}constant [1 x i64] [i64 99]
+// CK1: [[MTYPE06:@.+]] = {{.*}}constant [1 x i64] [i64 99]
+// CK1: [[MTYPE07:@.+]] = {{.*}}constant [1 x i64] [i64 99]
+// CK1: [[MTYPE08:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 35]
+// CK1: [[MTYPE09:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 99]
+// CK1: [[MTYPE10:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 99]
+// CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 35, i64 96]
+// CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 35, i64 96]
// CK1-LABEL: @_Z3foo
template<typename T>
@@ -42,7 +42,7 @@ void foo(float *&lr, T *&tr) {
T *t;
// CK1: [[T:%.+]] = load double*, double** [[DECL:@g]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+ // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to double**
// CK1: store double* [[T]], double** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE00]]
@@ -61,7 +61,7 @@ void foo(float *&lr, T *&tr) {
++g;
// CK1: [[T1:%.+]] = load float*, float** [[DECL:%.+]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+ // CK1: [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float**
// CK1: store float* [[T1]], float** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
@@ -92,7 +92,7 @@ void foo(float *&lr, T *&tr) {
++l;
// CK1: [[T1:%.+]] = load float*, float** [[DECL:%.+]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+ // CK1: [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float**
// CK1: store float* [[T1]], float** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
@@ -115,7 +115,7 @@ void foo(float *&lr, T *&tr) {
// CK1: [[BTHEN]]:
// CK1: [[T1:%.+]] = load float*, float** [[DECL:%.+]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+ // CK1: [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float**
// CK1: store float* [[T1]], float** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE04]]
@@ -152,7 +152,7 @@ void foo(float *&lr, T *&tr) {
// CK1: [[T2:%.+]] = load float**, float*** [[DECL:%.+]],
// CK1: [[T1:%.+]] = load float*, float** [[T2]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+ // CK1: [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float**
// CK1: store float* [[T1]], float** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE05]]
@@ -174,7 +174,7 @@ void foo(float *&lr, T *&tr) {
++lr;
// CK1: [[T1:%.+]] = load i32*, i32** [[DECL:%.+]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+ // CK1: [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
// CK1: store i32* [[T1]], i32** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE06]]
@@ -194,7 +194,7 @@ void foo(float *&lr, T *&tr) {
// CK1: [[T2:%.+]] = load i32**, i32*** [[DECL:%.+]],
// CK1: [[T1:%.+]] = load i32*, i32** [[T2]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
+ // CK1: [[BP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
// CK1: store i32* [[T1]], i32** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE07]]
@@ -216,7 +216,7 @@ void foo(float *&lr, T *&tr) {
++tr;
// CK1: [[T1:%.+]] = load float*, float** [[DECL:%.+]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32
+ // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float**
// CK1: store float* [[T1]], float** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE08]]
@@ -280,7 +280,7 @@ void foo(float *&lr, T *&tr) {
++l; ++t;
// CK1: [[T1:%.+]] = load i32*, i32** [[DECL:%.+]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0
+ // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
// CK1: store i32* [[T1]], i32** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE11]]
@@ -300,7 +300,7 @@ void foo(float *&lr, T *&tr) {
// CK1: [[T2:%.+]] = load i32**, i32*** [[DECL:%.+]],
// CK1: [[T1:%.+]] = load i32*, i32** [[T2]],
- // CK1: [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0
+ // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
// CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32**
// CK1: store i32* [[T1]], i32** [[CBP]],
// CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE12]]
@@ -348,7 +348,7 @@ void bar(float *&a, int *&b) {
// CK2: [[ST:%.+]] = type { double*, double** }
// CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
// CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
-// CK2: [[MTYPE02:@.+]] = {{.*}}constant [4 x i64] [i64 35, i64 19, i64 32, i64 844424930132048]
+// CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 35, i64 32, i64 562949953421392]
// CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 32, i64 281474976710739, i64 281474976710736]
template <typename T>
@@ -404,7 +404,7 @@ struct ST {
// CK2: getelementptr inbounds double, double* [[TTTT]], i32 1
b++;
- // CK2: [[BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32 3
+ // CK2: [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2
// CK2: [[CBP:%.+]] = bitcast i8** [[BP]] to double***
// CK2: store double** [[RVAL:%.+]], double*** [[CBP]],
// CK2: call void @__tgt_target_data_begin{{.+}}[[MTYPE02]]
diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp
index e63f19ad73c0..a394f5b1c3d5 100644
--- a/clang/test/OpenMP/target_map_codegen.cpp
+++ b/clang/test/OpenMP/target_map_codegen.cpp
@@ -3874,7 +3874,7 @@ int explicit_maps_template_args_and_members(int a){
// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE04:@.+]] = private {{.*}}constant [1 x i64] [i64 20]
-// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
@@ -3894,7 +3894,7 @@ int explicit_maps_template_args_and_members(int a){
// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 20]
-// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE10:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
@@ -3914,7 +3914,7 @@ int explicit_maps_template_args_and_members(int a){
// CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
// CK22: [[SIZE14:@.+]] = private {{.*}}constant [1 x i64] [i64 20]
-// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 51]
int a;
int c[100];
@@ -4010,11 +4010,10 @@ int explicit_maps_globals(void){
// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
- // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+ // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***
// CK22-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
- // CK22-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]]
+ // CK22-DAG: store i32** @d, i32*** [[CBP0]]
// CK22-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]]
- // CK22-DAG: [[RVAR0]] = load i32*, i32** @d
// CK22-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 2
// CK22-DAG: [[RVAR00]] = load i32*, i32** @d
@@ -4093,11 +4092,10 @@ int explicit_maps_globals(void){
// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
- // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+ // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]***
// CK22-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]**
- // CK22-DAG: store [[ST]]* [[RVAR0:%.+]], [[ST]]** [[CBP0]]
+ // CK22-DAG: store [[ST]]** @sd, [[ST]]*** [[CBP0]]
// CK22-DAG: store [[ST]]* [[SEC0:%.+]], [[ST]]** [[CP0]]
- // CK22-DAG: [[RVAR0]] = load [[ST]]*, [[ST]]** @sd
// CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[RVAR00:%.+]], i{{.+}} 2
// CK22-DAG: [[RVAR00]] = load [[ST]]*, [[ST]]** @sd
@@ -4176,11 +4174,10 @@ int explicit_maps_globals(void){
// CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
- // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[STT]]**
+ // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[STT]]***
// CK22-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[STT]]**
- // CK22-DAG: store [[STT]]* [[RVAR0:%.+]], [[STT]]** [[CBP0]]
+ // CK22-DAG: store [[STT]]** @std, [[STT]]*** [[CBP0]]
// CK22-DAG: store [[STT]]* [[SEC0:%.+]], [[STT]]** [[CP0]]
- // CK22-DAG: [[RVAR0]] = load [[STT]]*, [[STT]]** @std
// CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[STT]]* [[RVAR00:%.+]], i{{.+}} 2
// CK22-DAG: [[RVAR00]] = load [[STT]]*, [[STT]]** @std
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
index feea1129be4c..9eab8b3367c4 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -310,22 +310,23 @@ void device_side_scan(int arg) {
#ifdef CK5
-// CK5: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK5: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
// CK5-LABEL: lvalue
void lvalue(int *B, int l, int e) {
- // CK5-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK5-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
- // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK5-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
+ // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK5-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**
// CK5-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
- // CK5-DAG: store i32** [[B_ADDR:%.+]], i32*** [[BPC0]]
+ // CK5-DAG: store i32* [[B_VAL:%.+]], i32** [[BPC0]]
// CK5-DAG: store i32* [[B_VAL_2:%.+]], i32** [[PC0]]
+ // CK5-DAG: [[B_VAL]] = load i32*, i32** [[B_ADDR:%.+]]
// CK5-DAG: [[B_VAL_2]] = load i32*, i32** [[B_ADDR]]
#pragma omp target update to(*B)
*B += e;
@@ -351,28 +352,29 @@ void lvalue(int *B, int l, int e) {
#ifdef CK6
-// CK6: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK6: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK6: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
// CK6-LABEL: lvalue
void lvalue(int *B, int l, int e) {
- // CK6-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK6-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK6-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK6-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
- // CK6-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK6-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK6-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
+ // CK6-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK6-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK6-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**
// CK6-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
- // CK6-DAG: store i32** [[B_ADDR:%.+]], i32*** [[BPC0]]
+ // CK6-DAG: store i32* [[TWO:%.+]], i32** [[BPC0]]
// CK6-DAG: store i32* [[ADD_PTR:%.+]], i32** [[PC0]]
// CK6-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i{{32|64}} [[IDX_EXT:%.+]]
// CK6-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i{{32|64}} [[L_VAL:%.+]]
// CK6-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
// CK6-DAG: [[L_VAL]] = load i32, i32* [[L_ADDR:%.+]]
// CK6-DAG: store i32 {{.+}}, i32* [[L_ADDR]]
- // CK6-DAG: [[ONE]] = load i32*, i32** [[B_ADDR]]
+ // CK6-DAG: [[ONE]] = load i32*, i32** [[B_ADDR:%.+]]
+ // CK6-DAG: [[TWO]] = load i32*, i32** [[B_ADDR]]
#pragma omp target update to(*(B+l))
*(B+l) += e;
#pragma omp target update from(*(B+l))
@@ -397,25 +399,26 @@ void lvalue(int *B, int l, int e) {
#ifdef CK7
-// CK7: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK7: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK7: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
// CK7-LABEL: lvalue
void lvalue(int *B, int l, int e) {
- // CK7-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK7-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK7-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK7-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
- // CK7-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK7-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK7-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
+ // CK7-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK7-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK7-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**
// CK7-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
- // CK7-DAG: store i32** [[B_ADDR:%.+]], i32*** [[BPC0]]
+ // CK7-DAG: store i32* [[B_VAL:%.+]], i32** [[BPC0]]
// CK7-DAG: store i32* [[ARRAY_IDX:%.+]], i32** [[PC0]]
// CK7-DAG: [[ARRAY_IDX]] = getelementptr inbounds i32, i32* [[ADD_PTR:%.+]], i{{32|64}} [[IDX_PROM:%.+]]
// CK7-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i64 [[IDX_EXT:%.+]]
// CK7-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[B_VAL_2:%.+]], i32 [[L_VAL:%.+]]
+ // CK7-32-DAG: [[B_VAL]] = load i32*, i32** [[B_ADDR:%.+]]
// CK7-32-DAG: [[B_VAL_2]] = load i32*, i32** [[B_ADDR]]
// CK7-32-DAG: [[L_VAL]] = load i32, i32* [[L_ADDR:%.+]]
// CK7-32-DAG: [[IDX_PROM]] = load i32, i32* [[L_ADDR]]
@@ -446,18 +449,18 @@ void lvalue(int *B, int l, int e) {
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK8
-// CK8: [[SIZE00:@.+]] = {{.+}}constant [3 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK8: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 33, i64 16, i64 17]
+// CK8: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
+// CK8: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
// CK8-LABEL: lvalue
void lvalue(int **B, int l, int e) {
- // CK8-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}], [3 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK8-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}], [2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK8-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK8-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
- // CK8-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
- // CK8-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+ // CK8-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK8-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
// CK8-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
// CK8-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
// CK8-DAG: store i32** [[ARRAY_IDX_1:%.+]], i32*** [[BPC0]]
@@ -501,19 +504,19 @@ struct S {
double *p;
};
-// CK9: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673]
+// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
// CK9-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
- // CK9-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK9-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK9-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK9-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK9-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
//
- // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
- // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
- // CK9-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+ // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK9-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
// CK9-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
// CK9-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
// CK9-DAG: store double** [[P:%.+]], double*** [[BPC0]]
@@ -551,19 +554,19 @@ struct S {
double *p;
};
-// CK10: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673]
+// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
// CK10-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
- // CK10-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK10-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK10-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK10-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK10-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
//
- // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
- // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
- // CK10-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+ // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK10-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
// CK10-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
// CK10-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
// CK10-DAG: store double** [[P_VAL:%.+]], double*** [[BPC0]]
@@ -601,19 +604,19 @@ void lvalue(struct S *s, int l, int e) {
struct S {
double *p;
};
-// CK11: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673]
+// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
// CK11-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
- // CK11-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK11-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK11-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK11-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK11-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
//
- // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
- // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
- // CK11-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+ // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK11-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
// CK11-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
// CK11-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
// CK11-DAG: store double** [[P:%.+]], double*** [[BPC0]]
@@ -653,41 +656,44 @@ struct S {
double *p;
struct S *sp;
};
-// CK12: [[MTYPE00:@.+]] = {{.+}}constant [4 x i64] [i64 32, i64 281474976710657, i64 281474976710672, i64 17]
+// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710672, i64 17]
// CK12-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
- // CK12-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}, [4 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK12-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK12-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK12-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK12-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
//
- // CK12-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
- // CK12-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
- // CK12-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 3
+ // CK12-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+ // CK12-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+ // CK12-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
// CK12-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to double***
// CK12-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double**
// CK12-DAG: store double** [[P_VAL:%.+]], double*** [[BPC2]]
// CK12-DAG: store double* [[SIX:%.+]], double** [[PC2]]
// CK12-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE2]]
- // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
- // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
- // CK12-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+ // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+ // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+ // CK12-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
// CK12-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to [[STRUCT_S:%.+]]***
// CK12-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double***
// CK12-DAG: store [[STRUCT_S]]** [[SP:%.+]], [[STRUCT_S]]*** [[BPC1]]
// CK12-DAG: store double** [[P_VAL:%.+]], double*** [[PC1]]
// CK12-DAG: store i{{.+}} {{4|8}}, i{{.+}}* [[SIZE1]]
- // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK12-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+ // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK12-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0
// CK12-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[STRUCT_S:%.+]]**
// CK12-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to [[STRUCT_S]]***
+ // CK12-DAG: store [[STRUCT_S]]* [[ZERO:%.+]], [[STRUCT_S]]** [[BPC0]]
+ // CK12-DAG: store [[STRUCT_S]]** [[SP]], [[STRUCT_S]]*** [[PC0]]
// CK12-DAG: store [[STRUCT_S]]** [[S:%.+]], [[STRUCT_S]]*** [[S_VAL:%.+]]
// CK12-DAG: store i{{.+}} {{.+}}, i{{.+}}* [[SIZE0]]
// CK12-DAG: [[SP]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[ONE:%.+]], i32 0, i32 1
- // CK12-DAG: [[ONE]] = load %struct.S*, %struct.S** [[S]],
+ // CK12-DAG: [[ONE]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S:%.+]],
+ // CK12-DAG: [[ZERO]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S]],
#pragma omp target update to(*(s->sp->p))
*(s->sp->p) = e;
#pragma omp target update from(*(s->sp->p))
@@ -711,21 +717,21 @@ void lvalue(struct S *s, int l, int e) {
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK13
-// CK13: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 {{8|4}}, i64 4]
-// CK13: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK13: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK13: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
// CK13-LABEL: lvalue
void lvalue(int **BB, int a, int b) {
- // CK13-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK13-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK13-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK13-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
- // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK13-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32****
+ // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK13-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
// CK13-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
- // CK13-DAG: store i32*** [[BB_ADDR:%.+]], i32**** [[BPC0]]
+ // CK13-DAG: store i32** [[B_VAL1:%.+]], i32*** [[BPC0]]
// CK13-DAG: store i32* [[ADD_PTR_2:%.+]], i32** [[PC0]]
// CK13-64-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32, i32* [[RESULT:%.+]], i64 [[IDX_EXT_1:%.+]]
// CK13-32-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32, i32* [[RESULT:%.+]], i32 [[B_ADDR:%.+]]
@@ -734,7 +740,8 @@ void lvalue(int **BB, int a, int b) {
// CK13-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i64 [[IDX_EXT:%.+]]
// CK13-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i32 [[A_ADDR:%.+]]
// CK13-64-DAG: [[IDX_EXT]] = sext i32 [[TWO:%.+]] to i64
- // CK13-DAG: [[B_VAL]] = load i32**, i32*** [[BB_ADDR]]
+ // CK13-DAG: [[B_VAL]] = load i32**, i32*** [[BB_ADDR:%.+]]
+ // CK13-DAG: [[B_VAL1]] = load i32**, i32*** [[BB_ADDR]]
#pragma omp target update to(*(*(BB+a)+b))
*(*(BB+a)+b) = 1;
#pragma omp target update from(*(*(BB+a)+b))
@@ -831,7 +838,7 @@ void lvalue_member(SSA *sap) {
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK15
-// CK15: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673]
+// CK15: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
struct SSA {
double *p;
@@ -842,36 +849,27 @@ struct SSA {
//CK-15-LABEL: lvalue_member
void lvalue_member(SSA *sap) {
- // CK15-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK15-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK15-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK15-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK15-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
- // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
- // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
- // CK15-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
- // CK15-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to double***
- // CK15-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double**
- // CK15-DAG: store double** [[P_VAL:%.+]], double*** [[BPC2]]
- // CK15-DAG: store double* [[ADD_PTR:%.+]], double** [[PC2]]
- // CK15-DAG: store i64 8, i64* [[SIZE2]]
- // CK15-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} 3
- // CK15-DAG: [[THREE]] = load double*, double** [[P_VAL_1:%.+]]
- // CK15-DAG: [[P_VAL]] = getelementptr inbounds [[SSA:%.+]], [[SSA:%.+]]* [[THIS:%.+]], i32 0, i32 0
// CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
// CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
// CK15-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
- // CK15-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]**
- // CK15-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to [[SSA]]***
- // CK15-DAG: store [[SSA]]* [[SAP_VAL:%.+]], [[SSA]]** [[BPC1]],
- // CK15-DAG: store [[SSA]]** [[SAP_ADDR:%.+]], [[SSA]]*** [[PC1]]
- // CK15-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[SIZE1]]
- // CK15-DAG: [[SAP_VAL]] = load [[SSA]]*, [[SSA]]** [[SAP_ADDR]],
+ // CK15-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to double***
+ // CK15-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double**
+ // CK15-DAG: store double** [[P_VAL:%.+]], double*** [[BPC1]]
+ // CK15-DAG: store double* [[ADD_PTR:%.+]], double** [[PC1]]
+ // CK15-DAG: store i64 {{4|8}}, i64* [[SIZE1]]
+ // CK15-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} 3
+ // CK15-DAG: [[THREE]] = load double*, double** [[P_VAL_1:%.+]]
+ // CK15-DAG: [[P_VAL]] = getelementptr inbounds [[SSA:%.+]], [[SSA:%.+]]* [[THIS:%.+]], i32 0, i32 0
// CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK15-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0
- // CK15-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[SSA]]***
- // CK15-DAG: store [[SSA]]** [[SAP_ADDR]], [[SSA]]*** [[BPC0]],
+ // CK15-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[SSA]]**
+ // CK15-DAG: store [[SSA]]* [[ZERO:%.+]], [[SSA]]** [[BPC0]],
// CK15-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double***
// CK15-DAG: store double** [[P_VAL]], double*** [[PC0]],
// CK15-DAG: store i{{.+}} [[COMPUTE_SIZE:%.+]], i{{.+}}* [[SIZE0]]
@@ -881,6 +879,7 @@ void lvalue_member(SSA *sap) {
// CK15-DAG: [[EIGHT]] = ptrtoint i8* [[FIVE:%.+]] to i64
// CK15-DAG: [[SIX]] = bitcast double** {{.+}} to i8*
// CK15-DAG: [[FIVE]] = bitcast double** {{.+}} to i8*
+ // CK15-DAG: [[ZERO]] = load [[SSA]]*, [[SSA]]** %{{.+}},
#pragma omp target update to(*(3+sap->p))
*(3+sap->p) = 1;
#pragma omp target update from(*(3+sap->p))
@@ -904,25 +903,26 @@ void lvalue_member(SSA *sap) {
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK16
-// CK16: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 {{8|4}}, i64 4]
-// CK16: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK16: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK16: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
//CK16-LABEL: lvalue_find_base
void lvalue_find_base(float *f, int *i) {
- // CK16-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK16-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK16-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK16-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
- // CK16-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK16-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK16-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float***
+ // CK16-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK16-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK16-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
// CK16-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
- // CK16-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC0]]
+ // CK16-DAG: store float* [[F:%.+]], float** [[BPC0]]
// CK16-DAG: store float* [[ADD_PTR:%.+]], float** [[PC0]]
// CK16-32-DAG: [[ADD_PTR]] = getelementptr inbounds float, float* [[THREE:%.+]], i32 [[I:%.+]]
// CK16-64-DAG: [[ADD_PTR]] = getelementptr inbounds float, float* [[THREE:%.+]], i64 [[IDX_EXT:%.+]]
- // CK16-DAG: [[THREE]] = load float*, float** [[F_ADDR]],
+ // CK16-DAG: [[THREE]] = load float*, float** [[F_ADDR:%.+]],
+ // CK16-DAG: [[F]] = load float*, float** [[F_ADDR]],
// CK16-64-DAG: [[IDX_EXT]] = sext i32 [[I:%.+]] to i64
#pragma omp target update to(*(*i+f))
@@ -948,8 +948,8 @@ void lvalue_find_base(float *f, int *i) {
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK17
-// CK17: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 {{4|8}}, i64 4]
-// CK17: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK17: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK17: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
struct SSA {
int i;
@@ -959,15 +959,15 @@ struct SSA {
//CK17-LABEL: lvalue_find_base
void lvalue_find_base(float **f, SSA *sa) {
- // CK17-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK17-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK17-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK17-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
- // CK17-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK17-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK17-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float****
+ // CK17-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK17-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK17-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float***
// CK17-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
- // CK17-DAG: store float*** [[F_ADDR:%.+]], float**** [[BPC0]],
+ // CK17-DAG: store float** [[F_VAL:%.+]], float*** [[BPC0]],
// CK17-DAG: store float* [[ADD_PTR_4:%.+]], float** [[PC0]],
// CK17-64-DAG: [[ADD_PTR_4]] = getelementptr inbounds float, float* [[SEVEN:%.+]], i64 [[IDX_EXT_3:%.+]]
// CK17-64-DAG: [[IDX_EXT_3]] = sext i32 [[I_VAL:%.+]] to i64
@@ -981,6 +981,8 @@ void lvalue_find_base(float **f, SSA *sa) {
// CK17-DAG: [[FIVE]] = load i32, i32* [[I_2:%.+]],
// CK17-DAG: [[I_2]] = getelementptr inbounds [[SSA:%.+]], [[SSA]]* [[FOUR:%.+]], i32 0, i32 0
// CK17-DAG: [[FOUR]] = load [[SSA]]*, [[SSA]]** [[SSA_ADDR:%.+]],
+ // CK17-DAG: [[F]] = load float**, float*** [[F_ADDR:%.+]],
+ // CK17-DAG: [[F_VAL]] = load float**, float*** [[F_ADDR]],
#pragma omp target update to(*(sa->sa->i+*(1+sa->i+f)))
*(sa->sa->i+*(1+sa->i+f)) = 1;
@@ -1005,13 +1007,13 @@ void lvalue_find_base(float **f, SSA *sa) {
// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
#ifdef CK18
-// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 16]
-// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [2 x i64] [i64 34, i64 16]
+// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 34]
//CK18-LABEL: array_shaping
void array_shaping(float *f, int sa) {
- // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null)
+ // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null)
// CK18-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK18-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK18-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -1021,23 +1023,12 @@ void array_shaping(float *f, int sa) {
// CK18-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK18-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
- // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float***
+ // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
// CK18-DAG: store float* [[F1:%.+]], float** [[BPC0]],
- // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[PC0]],
- // CK18-DAG: store i64 {{8|4}}, i64* [[S0]],
- // CK18-DAG: [[F1]] = load float*, float** [[F_ADDR]],
-
- // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
-
- // CK18-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to float***
- // CK18-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to float**
-
- // CK18-DAG: store float** [[F_ADDR]], float*** [[BPC1]],
- // CK18-DAG: store float* [[F2:%.+]], float** [[PC1]],
- // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S1]],
+ // CK18-DAG: store float* [[F2:%.+]], float** [[PC0]],
+ // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
+ // CK18-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]],
// CK18-DAG: [[F2]] = load float*, float** [[F_ADDR]],
// CK18-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 4
@@ -1047,7 +1038,7 @@ void array_shaping(float *f, int sa) {
// CK18-32-DAG: [[SZ2]] = mul nuw i32 12, %{{.+}}
#pragma omp target update to(([3][sa][4])f)
sa = 1;
- // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null)
+ // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null)
// CK18-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK18-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK18-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -1057,23 +1048,12 @@ void array_shaping(float *f, int sa) {
// CK18-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK18-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
- // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float***
+ // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
// CK18-DAG: store float* [[F1:%.+]], float** [[BPC0]],
- // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[PC0]],
- // CK18-DAG: store i64 {{8|4}}, i64* [[S0]],
+ // CK18-DAG: store float* [[F2:%.+]], float** [[PC0]],
+ // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
// CK18-DAG: [[F1]] = load float*, float** [[F_ADDR]],
-
- // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
- // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
- // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
-
- // CK18-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to float***
- // CK18-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to float**
-
- // CK18-DAG: store float** [[F_ADDR]], float*** [[BPC1]],
- // CK18-DAG: store float* [[F2:%.+]], float** [[PC1]],
- // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S1]],
// CK18-DAG: [[F2]] = load float*, float** [[F_ADDR]],
// CK18-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 5
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 409089500015..9c7ab6309503 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -880,14 +880,9 @@ int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum, void **ArgBases,
return OFFLOAD_FAIL;
}
}
- } else if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
- TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBase, sizeof(void *), IsLast,
- false, IsHostPtr);
- TgtBaseOffset = 0; // no offset for ptrs.
- DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD " to "
- "object " DPxMOD "\n",
- DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBase), DPxPTR(HstPtrBase));
} else {
+ if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)
+ HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);
TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I], IsLast,
false, IsHostPtr);
TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
diff --git a/openmp/libomptarget/test/env/base_ptr_ref_count.c b/openmp/libomptarget/test/env/base_ptr_ref_count.c
new file mode 100644
index 000000000000..4036954182fa
--- /dev/null
+++ b/openmp/libomptarget/test/env/base_ptr_ref_count.c
@@ -0,0 +1,51 @@
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda
+// REQUIRES: libomptarget-debug
+
+#include <stdlib.h>
+
+int *allocate(size_t n) {
+ int *ptr = malloc(sizeof(int) * n);
+#pragma omp target enter data map(to : ptr[:n])
+ return ptr;
+}
+
+void deallocate(int *ptr, size_t n) {
+#pragma omp target exit data map(delete : ptr[:n])
+ free(ptr);
+}
+
+#pragma omp declare target
+int *cnt;
+void foo() {
+ ++(*cnt);
+}
+#pragma omp end declare target
+
+int main(void) {
+ int *A = allocate(10);
+ int *V = allocate(10);
+ deallocate(A, 10);
+ deallocate(V, 10);
+// CHECK-NOT: RefCount=2
+ cnt = malloc(sizeof(int));
+ *cnt = 0;
+#pragma omp target map(cnt[:1])
+ foo();
+ printf("Cnt = %d.\n", *cnt);
+// CHECK: Cnt = 1.
+ *cnt = 0;
+#pragma omp target data map(cnt[:1])
+#pragma omp target
+ foo();
+ printf("Cnt = %d.\n", *cnt);
+// CHECK: Cnt = 1.
+ free(cnt);
+
+ return 0;
+}
+
+
More information about the Openmp-commits
mailing list