| 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 | |
|---|
| 52 | int main() { |
|---|
| 53 | int arr[SIZE]; |
|---|
| 54 | |
|---|
| 55 | // CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]] |
|---|
| 56 | fprintf(stderr, format: "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, format: "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 map(present, tofrom : arr[SMALL]) |
|---|
| 67 | ; |
|---|
| 68 | } |
|---|
| 69 | |
|---|
| 70 | // CHECK: arr is present |
|---|
| 71 | fprintf(stderr, format: "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 error: Call to targetDataBegin failed, abort target. |
|---|
| 77 | // CHECK: omptarget error: Failed to process data before launching the kernel. |
|---|
| 78 | // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory |
|---|
| 79 | #pragma omp target data map(alloc : arr[SMALL]) |
|---|
| 80 | { |
|---|
| 81 | #pragma omp target map(present, tofrom : arr[LARGE]) |
|---|
| 82 | ; |
|---|
| 83 | } |
|---|
| 84 | |
|---|
| 85 | // CHECK-NOT: arr is present |
|---|
| 86 | fprintf(stderr, format: "arr is present\n" ); |
|---|
| 87 | |
|---|
| 88 | return 0; |
|---|
| 89 | } |
|---|
| 90 | |
|---|
