Skip to content

Commit 8e00703

Browse files
jyu2-gitjyu2-git
authored
[Clang][OpenMP] Fix runtime problem when explicit map both pointer and pointee (#92210)
ponter int *p for following map, test currently crash. map(p, p[:100]) or map(p, p[1]) Currly IR looks like // &p, &p, sizeof(int), TARGET_PARAM | TO | FROM // &p, p[0], 100sizeof(float) TO | FROM Worrking IR is // map(p, p[0:100]) to map(p[0:100]) // &p, &p[0], 100*sizeof(float), TARGET_PARAM | TO | FROM | PTR_AND_OBJ The change is add new argument AreBothBasePtrAndPteeMapped in generateInfoForComponentList Use that to skip map for map(p), when processing map(p[:100]) generate map with right flag.
1 parent dcf3102 commit 8e00703

File tree

3 files changed

+220
-9
lines changed

3 files changed

+220
-9
lines changed

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 28 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -6830,7 +6830,8 @@ class MappableExprsHandler {
68306830
const ValueDecl *Mapper = nullptr, bool ForDeviceAddr = false,
68316831
const ValueDecl *BaseDecl = nullptr, const Expr *MapExpr = nullptr,
68326832
ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
6833-
OverlappedElements = std::nullopt) const {
6833+
OverlappedElements = std::nullopt,
6834+
bool AreBothBasePtrAndPteeMapped = false) const {
68346835
// The following summarizes what has to be generated for each map and the
68356836
// types below. The generated information is expressed in this order:
68366837
// base pointer, section pointer, size, flags
@@ -7006,6 +7007,10 @@ class MappableExprsHandler {
70067007
// &(ps->p), &(ps->p[0]), 33*sizeof(double), MEMBER_OF(4) | PTR_AND_OBJ | TO
70077008
// (*) the struct this entry pertains to is the 4th element in the list
70087009
// of arguments, hence MEMBER_OF(4)
7010+
//
7011+
// map(p, p[:100])
7012+
// ===> map(p[:100])
7013+
// &p, &p[0], 100*sizeof(float), TARGET_PARAM | PTR_AND_OBJ | TO | FROM
70097014

70107015
// Track if the map information being generated is the first for a capture.
70117016
bool IsCaptureFirstInfo = IsFirstComponentList;
@@ -7029,6 +7034,8 @@ class MappableExprsHandler {
70297034
const auto *OASE = dyn_cast<ArraySectionExpr>(AssocExpr);
70307035
const auto *OAShE = dyn_cast<OMPArrayShapingExpr>(AssocExpr);
70317036

7037+
if (AreBothBasePtrAndPteeMapped && std::next(I) == CE)
7038+
return;
70327039
if (isa<MemberExpr>(AssocExpr)) {
70337040
// The base is the 'this' pointer. The content of the pointer is going
70347041
// to be the base of the field being mapped.
@@ -7071,8 +7078,9 @@ class MappableExprsHandler {
70717078
// can be associated with the combined storage if shared memory mode is
70727079
// active or the base declaration is not global variable.
70737080
const auto *VD = dyn_cast<VarDecl>(I->getAssociatedDeclaration());
7074-
if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
7075-
!VD || VD->hasLocalStorage())
7081+
if (!AreBothBasePtrAndPteeMapped &&
7082+
(CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
7083+
!VD || VD->hasLocalStorage()))
70767084
BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
70777085
else
70787086
FirstPointerInComplexData = true;
@@ -7394,11 +7402,13 @@ class MappableExprsHandler {
73947402
// same expression except for the first one. We also need to signal
73957403
// this map is the first one that relates with the current capture
73967404
// (there is a set of entries for each capture).
7397-
OpenMPOffloadMappingFlags Flags = getMapTypeBits(
7398-
MapType, MapModifiers, MotionModifiers, IsImplicit,
7399-
!IsExpressionFirstInfo || RequiresReference ||
7400-
FirstPointerInComplexData || IsMemberReference,
7401-
IsCaptureFirstInfo && !RequiresReference, IsNonContiguous);
7405+
OpenMPOffloadMappingFlags Flags =
7406+
getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit,
7407+
!IsExpressionFirstInfo || RequiresReference ||
7408+
FirstPointerInComplexData || IsMemberReference,
7409+
AreBothBasePtrAndPteeMapped ||
7410+
(IsCaptureFirstInfo && !RequiresReference),
7411+
IsNonContiguous);
74027412

74037413
if (!IsExpressionFirstInfo || IsMemberReference) {
74047414
// If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
@@ -8492,6 +8502,8 @@ class MappableExprsHandler {
84928502
assert(CurDir.is<const OMPExecutableDirective *>() &&
84938503
"Expect a executable directive");
84948504
const auto *CurExecDir = CurDir.get<const OMPExecutableDirective *>();
8505+
bool HasMapBasePtr = false;
8506+
bool HasMapArraySec = false;
84958507
for (const auto *C : CurExecDir->getClausesOfKind<OMPMapClause>()) {
84968508
const auto *EI = C->getVarRefs().begin();
84978509
for (const auto L : C->decl_component_lists(VD)) {
@@ -8503,6 +8515,11 @@ class MappableExprsHandler {
85038515
assert(VDecl == VD && "We got information for the wrong declaration??");
85048516
assert(!Components.empty() &&
85058517
"Not expecting declaration with no component lists.");
8518+
if (VD && E && VD->getType()->isAnyPointerType() && isa<DeclRefExpr>(E))
8519+
HasMapBasePtr = true;
8520+
if (VD && E && VD->getType()->isAnyPointerType() &&
8521+
(isa<ArraySectionExpr>(E) || isa<ArraySubscriptExpr>(E)))
8522+
HasMapArraySec = true;
85068523
DeclComponentLists.emplace_back(Components, C->getMapType(),
85078524
C->getMapTypeModifiers(),
85088525
C->isImplicit(), Mapper, E);
@@ -8685,7 +8702,9 @@ class MappableExprsHandler {
86858702
MapType, MapModifiers, std::nullopt, Components, CombinedInfo,
86868703
StructBaseCombinedInfo, PartialStruct, IsFirstComponentList,
86878704
IsImplicit, /*GenerateAllInfoForClauses*/ false, Mapper,
8688-
/*ForDeviceAddr=*/false, VD, VarRef);
8705+
/*ForDeviceAddr=*/false, VD, VarRef,
8706+
/*OverlappedElements*/ std::nullopt,
8707+
HasMapBasePtr && HasMapArraySec);
86898708
IsFirstComponentList = false;
86908709
}
86918710
}

‎clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp

Lines changed: 150 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,150 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
2+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
3+
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
4+
// RUN: %clang_cc1 -fopenmp -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
5+
6+
// expected-no-diagnostics
7+
#ifndef HEADER
8+
#define HEADER
9+
10+
extern void *malloc (int __size) throw () __attribute__ ((__malloc__));
11+
12+
void foo() {
13+
int *ptr = (int *) malloc(3 * sizeof(int));
14+
15+
#pragma omp target map(ptr, ptr[0:2])
16+
{
17+
ptr[1] = 6;
18+
}
19+
#pragma omp target map(ptr, ptr[2])
20+
{
21+
ptr[2] = 8;
22+
}
23+
}
24+
#endif
25+
// CHECK-LABEL: define {{[^@]+}}@_Z3foov
26+
// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
27+
// CHECK-NEXT: entry:
28+
// CHECK-NEXT: [[PTR:%.*]] = alloca ptr, align 8
29+
// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
30+
// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
31+
// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
32+
// CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
33+
// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [1 x ptr], align 8
34+
// CHECK-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [1 x ptr], align 8
35+
// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS4:%.*]] = alloca [1 x ptr], align 8
36+
// CHECK-NEXT: [[KERNEL_ARGS5:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS]], align 8
37+
// CHECK-NEXT: [[CALL:%.*]] = call noalias noundef ptr @_Z6malloci(i32 noundef signext 12) #[[ATTR3:[0-9]+]]
38+
// CHECK-NEXT: store ptr [[CALL]], ptr [[PTR]], align 8
39+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR]], align 8
40+
// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[PTR]], align 8
41+
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 0
42+
// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
43+
// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP2]], align 8
44+
// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
45+
// CHECK-NEXT: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
46+
// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
47+
// CHECK-NEXT: store ptr null, ptr [[TMP4]], align 8
48+
// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
49+
// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
50+
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
51+
// CHECK-NEXT: store i32 3, ptr [[TMP7]], align 4
52+
// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
53+
// CHECK-NEXT: store i32 1, ptr [[TMP8]], align 4
54+
// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
55+
// CHECK-NEXT: store ptr [[TMP5]], ptr [[TMP9]], align 8
56+
// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
57+
// CHECK-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8
58+
// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
59+
// CHECK-NEXT: store ptr @.offload_sizes, ptr [[TMP11]], align 8
60+
// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
61+
// CHECK-NEXT: store ptr @.offload_maptypes, ptr [[TMP12]], align 8
62+
// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
63+
// CHECK-NEXT: store ptr null, ptr [[TMP13]], align 8
64+
// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7
65+
// CHECK-NEXT: store ptr null, ptr [[TMP14]], align 8
66+
// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8
67+
// CHECK-NEXT: store i64 0, ptr [[TMP15]], align 8
68+
// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9
69+
// CHECK-NEXT: store i64 0, ptr [[TMP16]], align 8
70+
// CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10
71+
// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP17]], align 4
72+
// CHECK-NEXT: [[TMP18:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11
73+
// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP18]], align 4
74+
// CHECK-NEXT: [[TMP19:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12
75+
// CHECK-NEXT: store i32 0, ptr [[TMP19]], align 4
76+
// CHECK-NEXT: [[TMP20:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l15.region_id, ptr [[KERNEL_ARGS]])
77+
// CHECK-NEXT: [[TMP21:%.*]] = icmp ne i32 [[TMP20]], 0
78+
// CHECK-NEXT: br i1 [[TMP21]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]]
79+
// CHECK: omp_offload.failed:
80+
// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l15(ptr [[TMP0]]) #[[ATTR3]]
81+
// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]
82+
// CHECK: omp_offload.cont:
83+
// CHECK-NEXT: [[TMP22:%.*]] = load ptr, ptr [[PTR]], align 8
84+
// CHECK-NEXT: [[TMP23:%.*]] = load ptr, ptr [[PTR]], align 8
85+
// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP23]], i64 2
86+
// CHECK-NEXT: [[TMP24:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0
87+
// CHECK-NEXT: store ptr [[PTR]], ptr [[TMP24]], align 8
88+
// CHECK-NEXT: [[TMP25:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0
89+
// CHECK-NEXT: store ptr [[ARRAYIDX1]], ptr [[TMP25]], align 8
90+
// CHECK-NEXT: [[TMP26:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS4]], i64 0, i64 0
91+
// CHECK-NEXT: store ptr null, ptr [[TMP26]], align 8
92+
// CHECK-NEXT: [[TMP27:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i32 0, i32 0
93+
// CHECK-NEXT: [[TMP28:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i32 0, i32 0
94+
// CHECK-NEXT: [[TMP29:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 0
95+
// CHECK-NEXT: store i32 3, ptr [[TMP29]], align 4
96+
// CHECK-NEXT: [[TMP30:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 1
97+
// CHECK-NEXT: store i32 1, ptr [[TMP30]], align 4
98+
// CHECK-NEXT: [[TMP31:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 2
99+
// CHECK-NEXT: store ptr [[TMP27]], ptr [[TMP31]], align 8
100+
// CHECK-NEXT: [[TMP32:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 3
101+
// CHECK-NEXT: store ptr [[TMP28]], ptr [[TMP32]], align 8
102+
// CHECK-NEXT: [[TMP33:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 4
103+
// CHECK-NEXT: store ptr @.offload_sizes.1, ptr [[TMP33]], align 8
104+
// CHECK-NEXT: [[TMP34:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 5
105+
// CHECK-NEXT: store ptr @.offload_maptypes.2, ptr [[TMP34]], align 8
106+
// CHECK-NEXT: [[TMP35:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 6
107+
// CHECK-NEXT: store ptr null, ptr [[TMP35]], align 8
108+
// CHECK-NEXT: [[TMP36:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 7
109+
// CHECK-NEXT: store ptr null, ptr [[TMP36]], align 8
110+
// CHECK-NEXT: [[TMP37:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 8
111+
// CHECK-NEXT: store i64 0, ptr [[TMP37]], align 8
112+
// CHECK-NEXT: [[TMP38:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 9
113+
// CHECK-NEXT: store i64 0, ptr [[TMP38]], align 8
114+
// CHECK-NEXT: [[TMP39:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 10
115+
// CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP39]], align 4
116+
// CHECK-NEXT: [[TMP40:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 11
117+
// CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP40]], align 4
118+
// CHECK-NEXT: [[TMP41:%.*]] = getelementptr inbounds [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS5]], i32 0, i32 12
119+
// CHECK-NEXT: store i32 0, ptr [[TMP41]], align 4
120+
// CHECK-NEXT: [[TMP42:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l19.region_id, ptr [[KERNEL_ARGS5]])
121+
// CHECK-NEXT: [[TMP43:%.*]] = icmp ne i32 [[TMP42]], 0
122+
// CHECK-NEXT: br i1 [[TMP43]], label [[OMP_OFFLOAD_FAILED6:%.*]], label [[OMP_OFFLOAD_CONT7:%.*]]
123+
// CHECK: omp_offload.failed6:
124+
// CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l19(ptr [[TMP22]]) #[[ATTR3]]
125+
// CHECK-NEXT: br label [[OMP_OFFLOAD_CONT7]]
126+
// CHECK: omp_offload.cont7:
127+
// CHECK-NEXT: ret void
128+
//
129+
//
130+
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l15
131+
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR2:[0-9]+]] {
132+
// CHECK-NEXT: entry:
133+
// CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8
134+
// CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]], align 8
135+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
136+
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1
137+
// CHECK-NEXT: store i32 6, ptr [[ARRAYIDX]], align 4
138+
// CHECK-NEXT: ret void
139+
//
140+
//
141+
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l19
142+
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR2]] {
143+
// CHECK-NEXT: entry:
144+
// CHECK-NEXT: [[PTR_ADDR:%.*]] = alloca ptr, align 8
145+
// CHECK-NEXT: store ptr [[PTR]], ptr [[PTR_ADDR]], align 8
146+
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
147+
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
148+
// CHECK-NEXT: store i32 8, ptr [[ARRAYIDX]], align 4
149+
// CHECK-NEXT: ret void
150+
//

‎offload/test/mapping/map_both_pointer_pointee.c

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
2+
// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
3+
// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
4+
// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
5+
// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
6+
7+
// REQUIRES: unified_shared_memory
8+
// UNSUPPORTED: amdgcn-amd-amdhsa
9+
10+
#pragma omp declare target
11+
int *ptr1;
12+
#pragma omp end declare target
13+
14+
#include <stdio.h>
15+
#include <stdlib.h>
16+
int main() {
17+
ptr1 = (int *)malloc(sizeof(int) * 100);
18+
int *ptr2;
19+
ptr2 = (int *)malloc(sizeof(int) * 100);
20+
#pragma omp target map(ptr1, ptr1[ : 100])
21+
{ ptr1[1] = 6; }
22+
// CHECK: 6
23+
printf(" %d \n", ptr1[1]);
24+
#pragma omp target data map(ptr1[ : 5])
25+
{
26+
#pragma omp target map(ptr1[2], ptr1, ptr1[3]) map(ptr2, ptr2[2])
27+
{
28+
ptr1[2] = 7;
29+
ptr1[3] = 9;
30+
ptr2[2] = 7;
31+
}
32+
}
33+
// CHECK: 7 7 9
34+
printf(" %d %d %d \n", ptr2[2], ptr1[2], ptr1[3]);
35+
free(ptr1);
36+
#pragma omp target map(ptr2, ptr2[ : 100])
37+
{ ptr2[1] = 6; }
38+
// CHECK: 6
39+
printf(" %d \n", ptr2[1]);
40+
free(ptr2);
41+
return 0;
42+
}

0 commit comments

Comments
 (0)