| 1 | // RUN: %libomptarget-compile-and-run-generic |
| 2 | |
| 3 | #include <omp.h> |
| 4 | #include <stdlib.h> |
| 5 | |
| 6 | #define NUM_DIMS 3 |
| 7 | |
| 8 | int main() { |
| 9 | int d = omp_get_default_device(); |
| 10 | int id = omp_get_initial_device(); |
| 11 | int a[128], b[64], c[128], e[16], q[128], i; |
| 12 | void *p; |
| 13 | |
| 14 | if (d < 0 || d >= omp_get_num_devices()) |
| 15 | d = id; |
| 16 | |
| 17 | p = omp_target_alloc(130 * sizeof(int), d); |
| 18 | if (p == NULL) |
| 19 | return 0; |
| 20 | |
| 21 | for (i = 0; i < 128; i++) |
| 22 | q[i] = 0; |
| 23 | if (omp_target_memcpy(p, q, 128 * sizeof(int), 0, 0, d, id) != 0) |
| 24 | abort(); |
| 25 | |
| 26 | size_t volume[NUM_DIMS] = {2, 2, 3}; |
| 27 | size_t dst_offsets[NUM_DIMS] = {0, 0, 0}; |
| 28 | size_t src_offsets[NUM_DIMS] = {0, 0, 0}; |
| 29 | size_t dst_dimensions[NUM_DIMS] = {3, 4, 5}; |
| 30 | size_t src_dimensions[NUM_DIMS] = {2, 3, 4}; |
| 31 | |
| 32 | for (i = 0; i < 128; i++) |
| 33 | a[i] = 42; |
| 34 | for (i = 0; i < 64; i++) |
| 35 | b[i] = 24; |
| 36 | for (i = 0; i < 128; i++) |
| 37 | c[i] = 0; |
| 38 | for (i = 0; i < 16; i++) |
| 39 | e[i] = 77; |
| 40 | |
| 41 | omp_depend_t obj[2]; |
| 42 | |
| 43 | #pragma omp parallel num_threads(5) |
| 44 | #pragma omp single |
| 45 | { |
| 46 | #pragma omp task depend(out : p) |
| 47 | omp_target_memcpy(p, a, 128 * sizeof(int), 0, 0, d, id); |
| 48 | |
| 49 | #pragma omp task depend(inout : p) |
| 50 | omp_target_memcpy(p, b, 64 * sizeof(int), 0, 0, d, id); |
| 51 | |
| 52 | #pragma omp task depend(out : c) |
| 53 | for (i = 0; i < 128; i++) |
| 54 | c[i] = i + 1; |
| 55 | |
| 56 | #pragma omp depobj(obj[0]) depend(inout : p) |
| 57 | #pragma omp depobj(obj[1]) depend(in : c) |
| 58 | |
| 59 | /* This produces: 1 2 3 - - 5 6 7 - - at positions 0..9 and |
| 60 | 13 14 15 - - 17 18 19 - - at positions 20..29. */ |
| 61 | omp_target_memcpy_rect_async(p, c, sizeof(int), NUM_DIMS, volume, |
| 62 | dst_offsets, src_offsets, dst_dimensions, |
| 63 | src_dimensions, d, id, 2, obj); |
| 64 | |
| 65 | #pragma omp task depend(in : p) |
| 66 | omp_target_memcpy(p, e, 16 * sizeof(int), 0, 0, d, id); |
| 67 | } |
| 68 | |
| 69 | #pragma omp taskwait |
| 70 | |
| 71 | if (omp_target_memcpy(q, p, 128 * sizeof(int), 0, 0, id, d) != 0) |
| 72 | abort(); |
| 73 | |
| 74 | for (i = 0; i < 16; ++i) |
| 75 | if (q[i] != 77) |
| 76 | abort(); |
| 77 | if (q[20] != 13 || q[21] != 14 || q[22] != 15 || q[25] != 17 || q[26] != 18 || |
| 78 | q[27] != 19) |
| 79 | abort(); |
| 80 | for (i = 28; i < 64; ++i) |
| 81 | if (q[i] != 24) |
| 82 | abort(); |
| 83 | for (i = 64; i < 128; ++i) |
| 84 | if (q[i] != 42) |
| 85 | abort(); |
| 86 | |
| 87 | omp_target_free(p, d); |
| 88 | return 0; |
| 89 | } |
| 90 | |