Skip to content

Commit bb1cb6a

Browse files
abhinavgabaabhinavgaba
authored
[NFC][OpenMP] Add several use_device_ptr/addr tests. (#154939)
Most tests are either compfailing or runfailing. They should start passing once we start using ATTACH map-type based codegen. (#153683) Even after they start passing, there are a few places where the EXPECTED and actual CHECKs are different, due to two main issues: * use_device_ptr translation on `&p[0]` is not succeeding in looking-up a previously mapped `&p[1]` * privatization of byref use_device_addr operands is not happening correctly. The above should be fixed as separate standalone changes.
1 parent f0df62f commit bb1cb6a

15 files changed

+1481
-4
lines changed

‎offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_existing.cpp‎

Lines changed: 85 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,85 @@
1+
// RUN: %libomptarget-compilexx-run-and-check-generic
2+
3+
// XFAIL: *
4+
5+
#include <omp.h>
6+
#include <stdio.h>
7+
8+
// Test for various cases of use_device_addr on an array-section.
9+
// The corresponding data is mapped on a previous enter_data directive.
10+
11+
// Note that this tests for the current behavior wherein if a lookup fails,
12+
// the runtime returns nullptr, instead of the original host-address.
13+
// That was compatible with OpenMP 5.0, where it was a user error if
14+
// corresponding storage didn't exist, but with 5.1+, the runtime needs to
15+
// return the host address, as it needs to assume that the host-address is
16+
// device-accessible, as the user has guaranteed it.
17+
// Once the runtime returns the original host-address when the lookup fails, the
18+
// test will need to be updated.
19+
20+
int g, h[10];
21+
int *ph = &h[0];
22+
23+
struct S {
24+
int *paa[10][10];
25+
26+
void f1(int i) {
27+
paa[0][2] = &g;
28+
29+
int *original_ph3 = &ph[3];
30+
int **original_paa02 = &paa[0][2];
31+
32+
#pragma omp target enter data map(to : ph[3 : 4], paa[0][2 : 5])
33+
int *mapped_ptr_ph3 =
34+
(int *)omp_get_mapped_ptr(&ph[3], omp_get_default_device());
35+
int **mapped_ptr_paa02 =
36+
(int **)omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
37+
38+
// CHECK-COUNT-4: 1
39+
printf("%d\n", mapped_ptr_ph3 != nullptr);
40+
printf("%d\n", mapped_ptr_paa02 != nullptr);
41+
printf("%d\n", original_ph3 != mapped_ptr_ph3);
42+
printf("%d\n", original_paa02 != mapped_ptr_paa02);
43+
44+
// (A) use_device_addr operand within mapped address range.
45+
// CHECK: A: 1
46+
#pragma omp target data use_device_addr(ph[3 : 4])
47+
printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
48+
49+
// (B) use_device_addr operand in extended address range, but not
50+
// mapped address range.
51+
// CHECK: B: 1
52+
#pragma omp target data use_device_addr(ph[2])
53+
printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
54+
55+
// (C) use_device_addr/map: same base-array, different first-location.
56+
// CHECK: C: 1
57+
#pragma omp target data map(ph[3 : 2]) use_device_addr(ph[4 : 1])
58+
printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
59+
60+
// (D) use_device_addr/map: different base-array/pointers.
61+
// CHECK: D: 1
62+
#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
63+
printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
64+
65+
// (E) use_device_addr operand within mapped range of previous map.
66+
// CHECK: E: 1
67+
#pragma omp target data use_device_addr(paa[0])
68+
printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]);
69+
70+
// (F) use_device_addr/map: different operands, same base-array.
71+
// CHECK: F: 1
72+
#pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
73+
printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]);
74+
75+
// (G) use_device_addr/map: different base-array/pointers.
76+
// CHECK: G: 1
77+
#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
78+
printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]);
79+
80+
#pragma omp target exit data map(release : ph[3 : 4], paa[0][2 : 5])
81+
}
82+
};
83+
84+
S s1;
85+
int main() { s1.f1(1); }

‎offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp‎

Lines changed: 143 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,143 @@
1+
// RUN: %libomptarget-compilexx-run-and-check-generic
2+
3+
// XFAIL: *
4+
5+
#include <omp.h>
6+
#include <stdio.h>
7+
8+
// Test for various cases of use_device_addr on an array-section.
9+
// The corresponding data is not previously mapped.
10+
11+
// Note that this tests for the current behavior wherein if a lookup fails,
12+
// the runtime returns nullptr, instead of the original host-address.
13+
// That was compatible with OpenMP 5.0, where it was a user error if
14+
// corresponding storage didn't exist, but with 5.1+, the runtime needs to
15+
// return the host address, as it needs to assume that the host-address is
16+
// device-accessible, as the user has guaranteed it.
17+
// Once the runtime returns the original host-address when the lookup fails, the
18+
// test will need to be updated.
19+
20+
int g, h[10];
21+
int *ph = &h[0];
22+
23+
struct S {
24+
int *paa[10][10];
25+
26+
void f1(int i) {
27+
paa[0][2] = &g;
28+
29+
int *original_ph3 = &ph[3];
30+
int **original_paa02 = &paa[0][2];
31+
32+
// (A) No corresponding map, lookup should fail.
33+
// CHECK: A: 1 1 1
34+
#pragma omp target data use_device_addr(ph[3 : 4])
35+
{
36+
int *mapped_ptr_ph3 =
37+
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
38+
printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
39+
mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
40+
}
41+
42+
// (B) use_device_addr/map: different operands, same base-pointer.
43+
// use_device_addr operand within mapped address range.
44+
// CHECK: B: 1 1 1
45+
#pragma omp target data map(ph[2 : 3]) use_device_addr(ph[3 : 1])
46+
{
47+
int *mapped_ptr_ph4 =
48+
(int *)omp_get_mapped_ptr(original_ph3 + 1, omp_get_default_device());
49+
printf("B: %d %d %d\n", mapped_ptr_ph4 != nullptr,
50+
mapped_ptr_ph4 != original_ph3 + 1, &ph[4] == mapped_ptr_ph4);
51+
}
52+
53+
// (C) use_device_addr/map: different base-pointers.
54+
// No corresponding storage, lookup should fail.
55+
// CHECK: C: 1 1 1
56+
#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
57+
{
58+
int *mapped_ptr_ph3 =
59+
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
60+
printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
61+
mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
62+
}
63+
64+
// (D) use_device_addr/map: one of two maps with matching base-pointer.
65+
// use_device_addr operand within mapped address range of second map,
66+
// lookup should succeed.
67+
// CHECK: D: 1 1 1
68+
#pragma omp target data map(ph) map(ph[2 : 5]) use_device_addr(ph[3 : 4])
69+
{
70+
int *mapped_ptr_ph3 =
71+
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
72+
printf("D: %d %d %d\n", mapped_ptr_ph3 != nullptr,
73+
mapped_ptr_ph3 != original_ph3, &ph[3] == mapped_ptr_ph3);
74+
}
75+
76+
// (E) No corresponding map, lookup should fail
77+
// CHECK: E: 1 1 1
78+
#pragma omp target data use_device_addr(paa[0])
79+
{
80+
int **mapped_ptr_paa02 =
81+
(int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
82+
printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
83+
mapped_ptr_paa02 != original_paa02,
84+
&paa[0][2] == (int **)nullptr + 2);
85+
}
86+
87+
// (F) use_device_addr/map: different operands, same base-array.
88+
// use_device_addr within mapped address range. Lookup should succeed.
89+
// CHECK: F: 1 1 1
90+
#pragma omp target data map(paa) use_device_addr(paa[0])
91+
{
92+
int **mapped_ptr_paa02 =
93+
(int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
94+
printf("F: %d %d %d\n", mapped_ptr_paa02 != nullptr,
95+
mapped_ptr_paa02 != original_paa02,
96+
&paa[0][2] == mapped_ptr_paa02);
97+
}
98+
99+
// (G) use_device_addr/map: different operands, same base-array.
100+
// use_device_addr extends beyond existing mapping. Not spec compliant.
101+
// But the lookup succeeds because we use the base-address for translation.
102+
// CHECK: G: 1 1 1
103+
#pragma omp target data map(paa[0][4]) use_device_addr(paa[0])
104+
{
105+
int **mapped_ptr_paa04 = (int **)omp_get_mapped_ptr(
106+
original_paa02 + 2, omp_get_default_device());
107+
printf("G: %d %d %d\n", mapped_ptr_paa04 != nullptr,
108+
mapped_ptr_paa04 != original_paa02 + 2,
109+
&paa[0][4] == mapped_ptr_paa04);
110+
}
111+
112+
int *original_paa020 = &paa[0][2][0];
113+
int **original_paa0 = (int **)&paa[0];
114+
115+
// (H) use_device_addr/map: different base-pointers.
116+
// No corresponding storage for use_device_addr opnd, lookup should fail.
117+
// CHECK: H: 1 1 1
118+
#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0])
119+
{
120+
int **mapped_ptr_paa020 =
121+
(int **)omp_get_mapped_ptr(original_paa020, omp_get_default_device());
122+
int **mapped_ptr_paa0 =
123+
(int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
124+
printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
125+
mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
126+
}
127+
128+
// (I) use_device_addr/map: one map with different, one with same base-ptr.
129+
// Lookup should succeed.
130+
// CHECK: I: 1 1 1
131+
#pragma omp target data map(paa[0][2][0]) map(paa[0]) use_device_addr(paa[0][2])
132+
{
133+
int **mapped_ptr_paa02 =
134+
(int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
135+
printf("I: %d %d %d\n", mapped_ptr_paa02 != nullptr,
136+
mapped_ptr_paa02 != original_paa02,
137+
&paa[0][2] == mapped_ptr_paa02);
138+
}
139+
}
140+
};
141+
142+
S s1;
143+
int main() { s1.f1(1); }

‎offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_existing.cpp‎

Lines changed: 98 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,98 @@
1+
// RUN: %libomptarget-compilexx-run-and-check-generic
2+
3+
#include <omp.h>
4+
#include <stdio.h>
5+
6+
// Test for various cases of use_device_addr on an array-section on a reference.
7+
// The corresponding data is mapped on a previous enter_data directive.
8+
9+
// Note that this tests for the current behavior wherein if a lookup fails,
10+
// the runtime returns nullptr, instead of the original host-address.
11+
// That was compatible with OpenMP 5.0, where it was a user error if
12+
// corresponding storage didn't exist, but with 5.1+, the runtime needs to
13+
// return the host address, as it needs to assume that the host-address is
14+
// device-accessible, as the user has guaranteed it.
15+
// Once the runtime returns the original host-address when the lookup fails, the
16+
// test will need to be updated.
17+
18+
int g_ptee;
19+
int &g = g_ptee;
20+
21+
int h_ptee[10];
22+
int (&h)[10] = h_ptee;
23+
24+
int *ph_ptee = &h_ptee[0];
25+
int *&ph = ph_ptee;
26+
int *paa_ptee[10][10];
27+
28+
struct S {
29+
int *(&paa)[10][10] = paa_ptee;
30+
31+
void f1(int i) {
32+
paa[0][2] = &g;
33+
34+
int *original_ph3 = &ph[3];
35+
int **original_paa02 = &paa[0][2];
36+
37+
#pragma omp target enter data map(to : ph[3 : 4], paa[0][2 : 5])
38+
int *mapped_ptr_ph3 =
39+
(int *)omp_get_mapped_ptr(&ph[3], omp_get_default_device());
40+
int **mapped_ptr_paa02 =
41+
(int **)omp_get_mapped_ptr(&paa[0][2], omp_get_default_device());
42+
43+
// CHECK-COUNT-4: 1
44+
printf("%d\n", mapped_ptr_ph3 != nullptr);
45+
printf("%d\n", mapped_ptr_paa02 != nullptr);
46+
printf("%d\n", original_ph3 != mapped_ptr_ph3);
47+
printf("%d\n", original_paa02 != mapped_ptr_paa02);
48+
49+
// (A) use_device_addr operand within mapped address range.
50+
// EXPECTED: A: 1
51+
// CHECK: A: 0
52+
// FIXME: ph is not being privatized in the region.
53+
#pragma omp target data use_device_addr(ph[3 : 4])
54+
printf("A: %d\n", mapped_ptr_ph3 == &ph[3]);
55+
56+
// (B) use_device_addr operand in extended address range, but not
57+
// mapped address range.
58+
// EXPECTED: B: 1
59+
// CHECK: B: 0
60+
// FIXME: ph is not being privatized in the region.
61+
#pragma omp target data use_device_addr(ph[2])
62+
printf("B: %d\n", mapped_ptr_ph3 == &ph[3]);
63+
64+
// (C) use_device_addr/map: same base-array, different first-location.
65+
// EXPECTED: C: 1
66+
// CHECK: C: 0
67+
// FIXME: ph is not being privatized in the region.
68+
#pragma omp target data map(ph[3 : 2]) use_device_addr(ph[4 : 1])
69+
printf("C: %d\n", mapped_ptr_ph3 == &ph[3]);
70+
71+
// (D) use_device_addr/map: different base-array/pointers.
72+
// EXPECTED: D: 1
73+
// CHECK: D: 0
74+
// FIXME: ph is not being privatized in the region.
75+
#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
76+
printf("D: %d\n", mapped_ptr_ph3 == &ph[3]);
77+
78+
// (E) use_device_addr operand within mapped range of previous map.
79+
// CHECK: E: 1
80+
#pragma omp target data use_device_addr(paa[0])
81+
printf("E: %d\n", mapped_ptr_paa02 == &paa[0][2]);
82+
83+
// (F) use_device_addr/map: different operands, same base-array.
84+
// CHECK: F: 1
85+
#pragma omp target data map(paa[0][3]) use_device_addr(paa[0][2])
86+
printf("F: %d\n", mapped_ptr_paa02 == &paa[0][2]);
87+
88+
// (G) use_device_addr/map: different base-array/pointers.
89+
// CHECK: G: 1
90+
#pragma omp target data map(paa[0][2][0]) use_device_addr(paa[0][2])
91+
printf("G: %d\n", mapped_ptr_paa02 == &paa[0][2]);
92+
93+
#pragma omp target exit data map(release : ph[3 : 4], paa[0][2 : 5])
94+
}
95+
};
96+
97+
S s1;
98+
int main() { s1.f1(1); }

0 commit comments

Comments
 (0)