1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
|
// --------------------------------------------------
// Check extends before
// --------------------------------------------------
// RUN: %libomptarget-compile-generic \
// RUN: -fopenmp-version=51 -DEXTENDS=BEFORE
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
// --------------------------------------------------
// Check extends after
// --------------------------------------------------
// RUN: %libomptarget-compile-generic \
// RUN: -fopenmp-version=51 -DEXTENDS=AFTER
// RUN: %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
// END.
#include <stdio.h>
#define BEFORE 0
#define AFTER 1
#define SIZE 100
#if EXTENDS == BEFORE
# define SMALL_BEG (SIZE-2)
# define SMALL_END SIZE
# define LARGE_BEG 0
# define LARGE_END SIZE
#elif EXTENDS == AFTER
# define SMALL_BEG 0
# define SMALL_END 2
# define LARGE_BEG 0
# define LARGE_END SIZE
#else
# error EXTENDS undefined
#endif
#define SMALL SMALL_BEG:(SMALL_END-SMALL_BEG)
#define LARGE LARGE_BEG:(LARGE_END-LARGE_BEG)
void check_not_present() {
int arr[SIZE];
for (int i = 0; i < SIZE; ++i)
arr[i] = 99;
// CHECK-LABEL: checking not present
fprintf(stderr, "checking not present\n");
// arr[LARGE] isn't (fully) present at the end of the target data region, so
// the device-to-host transfer should not be performed, or it might fail.
#pragma omp target data map(tofrom: arr[LARGE])
{
#pragma omp target exit data map(delete: arr[LARGE])
#pragma omp target enter data map(alloc: arr[SMALL])
#pragma omp target map(alloc: arr[SMALL])
for (int i = SMALL_BEG; i < SMALL_END; ++i)
arr[i] = 88;
}
// CHECK-NOT: Libomptarget
// CHECK-NOT: error
for (int i = 0; i < SIZE; ++i) {
if (arr[i] != 99)
fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
}
}
void check_is_present() {
int arr[SIZE];
for (int i = 0; i < SIZE; ++i)
arr[i] = 99;
// CHECK-LABEL: checking is present
fprintf(stderr, "checking is present\n");
// arr[SMALL] is (fully) present at the end of the target data region, and the
// device-to-host transfer should be performed only for it even though more
// of the array is then present.
#pragma omp target data map(tofrom: arr[SMALL])
{
#pragma omp target exit data map(delete: arr[SMALL])
#pragma omp target enter data map(alloc: arr[LARGE])
#pragma omp target map(alloc: arr[LARGE])
for (int i = LARGE_BEG; i < LARGE_END; ++i)
arr[i] = 88;
}
// CHECK-NOT: Libomptarget
// CHECK-NOT: error
for (int i = 0; i < SIZE; ++i) {
if (SMALL_BEG <= i && i < SMALL_END) {
if (arr[i] != 88)
fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
} else if (arr[i] != 99) {
fprintf(stderr, "error: arr[%d]=%d\n", i, arr[i]);
}
}
}
int main() {
check_not_present();
check_is_present();
return 0;
}
|
