xref: /llvm-project/offload/test/api/omp_target_memcpy_async2.c (revision 330d8983d25d08580fc1642fea48b2473f47a9da)
1 // RUN: %libomptarget-compile-and-run-generic
2 
3 #include "stdio.h"
4 #include <omp.h>
5 #include <stdlib.h>
6 
main()7 int main() {
8   int d = omp_get_default_device();
9   int id = omp_get_initial_device();
10   int a[128], b[64], c[32], e[16], q[128], i;
11   void *p;
12 
13   if (d < 0 || d >= omp_get_num_devices())
14     d = id;
15 
16   p = omp_target_alloc(130 * sizeof(int), d);
17   if (p == NULL)
18     return 0;
19 
20   for (i = 0; i < 128; ++i)
21     a[i] = i + 1;
22   for (i = 0; i < 64; ++i)
23     b[i] = i + 2;
24   for (i = 0; i < 32; i++)
25     c[i] = 0;
26   for (i = 0; i < 16; i++)
27     e[i] = i + 4;
28 
29   omp_depend_t obj[2];
30 
31 #pragma omp parallel num_threads(5)
32 #pragma omp single
33   {
34 #pragma omp task depend(out : p)
35     omp_target_memcpy(p, a, 128 * sizeof(int), 0, 0, d, id);
36 
37 #pragma omp task depend(inout : p)
38     omp_target_memcpy(p, b, 64 * sizeof(int), 0, 0, d, id);
39 
40 #pragma omp task depend(out : c)
41     for (i = 0; i < 32; i++)
42       c[i] = i + 3;
43 
44 #pragma omp depobj(obj[0]) depend(inout : p)
45 #pragma omp depobj(obj[1]) depend(in : c)
46     omp_target_memcpy_async(p, c, 32 * sizeof(int), 0, 0, d, id, 2, obj);
47 
48 #pragma omp task depend(in : p)
49     omp_target_memcpy(p, e, 16 * sizeof(int), 0, 0, d, id);
50   }
51 
52 #pragma omp taskwait
53 
54   for (i = 0; i < 128; ++i)
55     q[i] = 0;
56   omp_target_memcpy(q, p, 128 * sizeof(int), 0, 0, id, d);
57   for (i = 0; i < 16; ++i)
58     if (q[i] != i + 4)
59       abort();
60   for (i = 16; i < 32; ++i)
61     if (q[i] != i + 3)
62       abort();
63   for (i = 32; i < 64; ++i)
64     if (q[i] != i + 2)
65       abort();
66   for (i = 64; i < 128; ++i)
67     if (q[i] != i + 1)
68       abort();
69 
70   omp_target_free(p, d);
71 
72   return 0;
73 }
74