1 // --------------------------------------------------
2 // Check extends before
3 // --------------------------------------------------
4
5 // RUN: %libomptarget-compile-generic \
6 // RUN: -DEXTENDS=BEFORE
7 // RUN: %libomptarget-run-fail-generic 2>&1 \
8 // RUN: | %fcheck-generic
9
10 // --------------------------------------------------
11 // Check extends after
12 // --------------------------------------------------
13
14 // RUN: %libomptarget-compile-generic \
15 // RUN: -DEXTENDS=AFTER
16 // RUN: %libomptarget-run-fail-generic 2>&1 \
17 // RUN: | %fcheck-generic
18
19 // END.
20
21 #include <stdio.h>
22
23 #define BEFORE 0
24 #define AFTER 1
25
26 #define SIZE 100
27
28 #if EXTENDS == BEFORE
29 #define SMALL_BEG (SIZE - 2)
30 #define SMALL_END SIZE
31 #define LARGE_BEG 0
32 #define LARGE_END SIZE
33 #elif EXTENDS == AFTER
34 #define SMALL_BEG 0
35 #define SMALL_END 2
36 #define LARGE_BEG 0
37 #define LARGE_END SIZE
38 #else
39 #error EXTENDS undefined
40 #endif
41
42 #define SMALL_SIZE (SMALL_END - SMALL_BEG)
43 #define LARGE_SIZE (LARGE_END - LARGE_BEG)
44
45 #define SMALL \
46 SMALL_BEG: \
47 SMALL_SIZE
48 #define LARGE \
49 LARGE_BEG: \
50 LARGE_SIZE
51
main()52 int main() {
53 int arr[SIZE];
54
55 // CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]]
56 fprintf(stderr, "addr=%p, size=%ld\n", &arr[SMALL_BEG],
57 SMALL_SIZE * sizeof arr[0]);
58
59 // CHECK: addr=0x[[#%x,LARGE_ADDR:]], size=[[#%u,LARGE_BYTES:]]
60 fprintf(stderr, "addr=%p, size=%ld\n", &arr[LARGE_BEG],
61 LARGE_SIZE * sizeof arr[0]);
62
63 // CHECK-NOT: omptarget
64 #pragma omp target data map(alloc : arr[LARGE])
65 {
66 #pragma omp target data map(present, tofrom : arr[SMALL])
67 ;
68 }
69
70 // CHECK: arr is present
71 fprintf(stderr, "arr is present\n");
72
73 // CHECK: omptarget 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)
74 // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes)
75 // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier).
76 // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory
77 #pragma omp target data map(alloc : arr[SMALL])
78 {
79 #pragma omp target data map(present, tofrom : arr[LARGE])
80 ;
81 }
82
83 // CHECK-NOT: arr is present
84 fprintf(stderr, "arr is present\n");
85
86 return 0;
87 }
88