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
|
// --------------------------------------------------
// Check extends before
// --------------------------------------------------
// RUN: %libomptarget-compile-generic \
// RUN: -DEXTENDS=BEFORE
// RUN: %libomptarget-run-fail-generic 2>&1 \
// RUN: | %fcheck-generic
// --------------------------------------------------
// Check extends after
// --------------------------------------------------
// RUN: %libomptarget-compile-generic \
// RUN: -DEXTENDS=AFTER
// RUN: %libomptarget-run-fail-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_SIZE (SMALL_END - SMALL_BEG)
#define LARGE_SIZE (LARGE_END - LARGE_BEG)
#define SMALL \
SMALL_BEG: \
SMALL_SIZE
#define LARGE \
LARGE_BEG: \
LARGE_SIZE
int main() {
int arr[SIZE];
// CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]]
fprintf(stderr, "addr=%p, size=%ld\n", &arr[SMALL_BEG],
SMALL_SIZE * sizeof arr[0]);
// CHECK: addr=0x[[#%x,LARGE_ADDR:]], size=[[#%u,LARGE_BYTES:]]
fprintf(stderr, "addr=%p, size=%ld\n", &arr[LARGE_BEG],
LARGE_SIZE * sizeof arr[0]);
// CHECK-NOT: Libomptarget
#pragma omp target data map(alloc : arr[LARGE])
{
#pragma omp target map(present, tofrom : arr[SMALL])
;
}
// CHECK: arr is present
fprintf(stderr, "arr is present\n");
// CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes)
// CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes)
// CHECK: Libomptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
// CHECK: Libomptarget error: Call to targetDataBegin failed, abort target.
// CHECK: Libomptarget error: Failed to process data before launching the kernel.
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
#pragma omp target data map(alloc : arr[SMALL])
{
#pragma omp target map(present, tofrom : arr[LARGE])
;
}
// CHECK-NOT: arr is present
fprintf(stderr, "arr is present\n");
return 0;
}
|
