1283514Sarybchikextern
2283514Sarybchik#ifdef __cplusplus
3283514Sarybchik"C"
4283514Sarybchik#endif
5283514Sarybchikvoid abort (void);
6283514Sarybchik
7283514Sarybchikvoid
8283514Sarybchikfn1 (double *x, double *y, int z)
9283514Sarybchik{
10283514Sarybchik  int i;
11283514Sarybchik  for (i = 0; i < z; i++)
12283514Sarybchik    {
13283514Sarybchik      x[i] = i & 31;
14283514Sarybchik      y[i] = (i & 63) - 30;
15283514Sarybchik    }
16283514Sarybchik}
17283514Sarybchik
18283514Sarybchik#pragma omp declare target
19283514Sarybchikint tgtv = 6;
20283514Sarybchikint
21283514Sarybchiktgt (void)
22283514Sarybchik{
23283514Sarybchik  #pragma omp atomic update
24283514Sarybchik    tgtv++;
25283514Sarybchik  return 0;
26283514Sarybchik}
27283514Sarybchik#pragma omp end declare target
28283514Sarybchik
29283514Sarybchikdouble
30283514Sarybchikfn2 (int x, int y, int z)
31283514Sarybchik{
32283514Sarybchik  double b[1024], c[1024], s = 0;
33283514Sarybchik  int i, j;
34283514Sarybchik  fn1 (b, c, x);
35283514Sarybchik  #pragma omp target data map(to: b)
36283514Sarybchik  {
37283514Sarybchik    #pragma omp target map(tofrom: c)
38283514Sarybchik      #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x)
39283514Sarybchik	#pragma omp distribute dist_schedule(static, 4) collapse(1)
40283514Sarybchik	  for (j=0; j < x; j += y)
41283514Sarybchik	    #pragma omp parallel for reduction(+:s)
42283514Sarybchik	      for (i = j; i < j + y; i++)
43283514Sarybchik		tgt (), s += b[i] * c[i];
44283514Sarybchik    #pragma omp target update from(b, tgtv)
45283514Sarybchik  }
46283514Sarybchik  return s;
47283514Sarybchik}
48283514Sarybchik
49283514Sarybchikdouble
50283514Sarybchikfn3 (int x)
51283514Sarybchik{
52283514Sarybchik  double b[1024], c[1024], s = 0;
53283514Sarybchik  int i;
54283514Sarybchik  fn1 (b, c, x);
55283514Sarybchik  #pragma omp target map(to: b, c)
56283514Sarybchik    #pragma omp parallel for reduction(+:s)
57283514Sarybchik      for (i = 0; i < x; i++)
58283514Sarybchik	tgt (), s += b[i] * c[i];
59283514Sarybchik  return s;
60283514Sarybchik}
61283514Sarybchik
62283514Sarybchikdouble
63283514Sarybchikfn4 (int x, double *p)
64283514Sarybchik{
65283514Sarybchik  double b[1024], c[1024], d[1024], s = 0;
66283514Sarybchik  int i;
67283514Sarybchik  fn1 (b, c, x);
68283514Sarybchik  fn1 (d + x, p + x, x);
69283514Sarybchik  #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)])
70283514Sarybchik    #pragma omp parallel for reduction(+:s)
71283514Sarybchik      for (i = 0; i < x; i++)
72283514Sarybchik	s += b[i] * c[i] + d[x + i] + p[x + i];
73283514Sarybchik  return s;
74283514Sarybchik}
75283514Sarybchik
76283514Sarybchikint
77283514Sarybchikmain ()
78283514Sarybchik{
79283514Sarybchik  double a = fn2 (128, 4, 6);
80283514Sarybchik  int b = tgtv;
81283514Sarybchik  double c = fn3 (61);
82283514Sarybchik  #pragma omp target update from(tgtv)
83283514Sarybchik  int d = tgtv;
84283514Sarybchik  double e[1024];
85283514Sarybchik  double f = fn4 (64, e);
86283514Sarybchik  if (a != 13888.0 || b != 6 + 128 || c != 4062.0 || d != 6 + 128 + 61
87283514Sarybchik      || f != 8032.0)
88283514Sarybchik    abort ();
89283514Sarybchik  return 0;
90283514Sarybchik}
91283514Sarybchik