blob: 5398905f4113bf7c9a72a53d89542ddfd08c0591 [file] [log] [blame]
#include <openacc.h>
#include <stdlib.h>
#include <stdbool.h>
#define N 32
int
main(int argc, char **argv)
{
float *a, *b, *d_a, *d_b, exp, exp2;
int i;
const int one = 1;
const int zero = 0;
int n;
a = (float *) malloc (N * sizeof (float));
b = (float *) malloc (N * sizeof (float));
d_a = (float *) acc_malloc (N * sizeof (float));
d_b = (float *) acc_malloc (N * sizeof (float));
for (i = 0; i < N; i++)
a[i] = 4.0;
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(1)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 5.0;
#else
exp = 4.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 16.0;
#pragma acc parallel if(0)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 17.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 8.0;
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(one)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 9.0;
#else
exp = 8.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 22.0;
#pragma acc parallel if(zero)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 23.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 16.0;
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(true)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 17.0;
#else
exp = 16.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 76.0;
#pragma acc parallel if(false)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 77.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 22.0;
n = 1;
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(n)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 23.0;
#else
exp = 22.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 18.0;
n = 0;
#pragma acc parallel if(n)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 19.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 49.0;
n = 1;
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(n + n)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 50.0;
#else
exp = 49.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 38.0;
n = 0;
#pragma acc parallel if(n + n)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 39.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 91.0;
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(-2)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 92.0;
#else
exp = 91.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 43.0;
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(one == 1)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 44.0;
#else
exp = 43.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 87.0;
#pragma acc parallel if(one == 0)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 88.0)
abort();
}
for (i = 0; i < N; i++)
{
a[i] = 3.0;
b[i] = 9.0;
}
#if ACC_MEM_SHARED
exp = 0.0;
exp2 = 0.0;
#else
acc_map_data (a, d_a, N * sizeof (float));
acc_map_data (b, d_b, N * sizeof (float));
exp = 3.0;
exp2 = 9.0;
#endif
#pragma acc update device(a[0:N], b[0:N]) if(1)
for (i = 0; i < N; i++)
{
a[i] = 0.0;
b[i] = 0.0;
}
#pragma acc update host(a[0:N], b[0:N]) if(1)
for (i = 0; i < N; i++)
{
if (a[i] != exp)
abort();
if (b[i] != exp2)
abort();
}
for (i = 0; i < N; i++)
{
a[i] = 6.0;
b[i] = 12.0;
}
#pragma acc update device(a[0:N], b[0:N]) if(0)
for (i = 0; i < N; i++)
{
a[i] = 0.0;
b[i] = 0.0;
}
#pragma acc update host(a[0:N], b[0:N]) if(1)
for (i = 0; i < N; i++)
{
if (a[i] != exp)
abort();
if (b[i] != exp2)
abort();
}
for (i = 0; i < N; i++)
{
a[i] = 26.0;
b[i] = 21.0;
}
#pragma acc update device(a[0:N], b[0:N]) if(1)
for (i = 0; i < N; i++)
{
a[i] = 0.0;
b[i] = 0.0;
}
#pragma acc update host(a[0:N], b[0:N]) if(0)
for (i = 0; i < N; i++)
{
if (a[i] != 0.0)
abort();
if (b[i] != 0.0)
abort();
}
#if !ACC_MEM_SHARED
acc_unmap_data (a);
acc_unmap_data (b);
#endif
acc_free (d_a);
acc_free (d_b);
for (i = 0; i < N; i++)
{
a[i] = 4.0;
b[i] = 0.0;
}
#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(1)
{
#pragma acc parallel present(a[0:N])
{
int ii;
for (ii = 0; ii < N; ii++)
{
b[ii] = a[ii];
}
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 4.0)
abort();
}
for (i = 0; i < N; i++)
{
a[i] = 8.0;
b[i] = 1.0;
}
#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(0)
{
#if !ACC_MEM_SHARED
if (acc_is_present (a, N * sizeof (float)))
abort ();
#endif
#if !ACC_MEM_SHARED
if (acc_is_present (b, N * sizeof (float)))
abort ();
#endif
}
for (i = 0; i < N; i++)
{
a[i] = 18.0;
b[i] = 21.0;
}
#pragma acc data copyin(a[0:N]) if(1)
{
#if !ACC_MEM_SHARED
if (!acc_is_present (a, N * sizeof (float)))
abort ();
#endif
#pragma acc data copyout(b[0:N]) if(0)
{
#if !ACC_MEM_SHARED
if (acc_is_present (b, N * sizeof (float)))
abort ();
#endif
#pragma acc data copyout(b[0:N]) if(1)
{
#pragma acc parallel present(a[0:N]) present(b[0:N])
{
int ii;
for (ii = 0; ii < N; ii++)
{
b[ii] = a[ii];
}
}
}
#if !ACC_MEM_SHARED
if (acc_is_present (b, N * sizeof (float)))
abort ();
#endif
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 18.0)
abort ();
}
#pragma acc enter data copyin (b[0:N]) if (0)
#if !ACC_MEM_SHARED
if (acc_is_present (b, N * sizeof (float)))
abort ();
#endif
#pragma acc exit data delete (b[0:N]) if (0)
#pragma acc enter data copyin (b[0:N]) if (1)
#if !ACC_MEM_SHARED
if (!acc_is_present (b, N * sizeof (float)))
abort ();
#endif
#pragma acc exit data delete (b[0:N]) if (1)
#if !ACC_MEM_SHARED
if (acc_is_present (b, N * sizeof (float)))
abort ();
#endif
#pragma acc enter data copyin (b[0:N]) if (zero)
#if !ACC_MEM_SHARED
if (acc_is_present (b, N * sizeof (float)))
abort ();
#endif
#pragma acc exit data delete (b[0:N]) if (zero)
#pragma acc enter data copyin (b[0:N]) if (one)
#if !ACC_MEM_SHARED
if (!acc_is_present (b, N * sizeof (float)))
abort ();
#endif
#pragma acc exit data delete (b[0:N]) if (one)
#if !ACC_MEM_SHARED
if (acc_is_present (b, N * sizeof (float)))
abort ();
#endif
#pragma acc enter data copyin (b[0:N]) if (one == 0)
#if !ACC_MEM_SHARED
if (acc_is_present (b, N * sizeof (float)))
abort ();
#endif
#pragma acc exit data delete (b[0:N]) if (one == 0)
#pragma acc enter data copyin (b[0:N]) if (one == 1)
#if !ACC_MEM_SHARED
if (!acc_is_present (b, N * sizeof (float)))
abort ();
#endif
#pragma acc exit data delete (b[0:N]) if (one == 1)
#if !ACC_MEM_SHARED
if (acc_is_present (b, N * sizeof (float)))
abort ();
#endif
for (i = 0; i < N; i++)
a[i] = 4.0;
#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(1)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 5.0;
#else
exp = 4.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 16.0;
#pragma acc kernels if(0)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 17.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 8.0;
#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(one)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 9.0;
#else
exp = 8.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 22.0;
#pragma acc kernels if(zero)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 23.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 16.0;
#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(true)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 17.0;
#else
exp = 16.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 76.0;
#pragma acc kernels if(false)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 77.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 22.0;
n = 1;
#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(n)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 23.0;
#else
exp = 22.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 18.0;
n = 0;
#pragma acc kernels if(n)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 19.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 49.0;
n = 1;
#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(n + n)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 50.0;
#else
exp = 49.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 38.0;
n = 0;
#pragma acc kernels if(n + n)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 39.0)
abort();
}
for (i = 0; i < N; i++)
a[i] = 91.0;
#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(-2)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 92.0;
#else
exp = 91.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 43.0;
#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) if(one == 1)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
#if ACC_MEM_SHARED
exp = 44.0;
#else
exp = 43.0;
#endif
for (i = 0; i < N; i++)
{
if (b[i] != exp)
abort();
}
for (i = 0; i < N; i++)
a[i] = 87.0;
#pragma acc kernels if(one == 0)
{
int ii;
for (ii = 0; ii < N; ii++)
{
if (acc_on_device (acc_device_host))
b[ii] = a[ii] + 1;
else
b[ii] = a[ii];
}
}
for (i = 0; i < N; i++)
{
if (b[i] != 88.0)
abort();
}
for (i = 0; i < N; i++)
{
a[i] = 3.0;
b[i] = 9.0;
}
#if ACC_MEM_SHARED
exp = 0.0;
exp2 = 0.0;
#else
acc_map_data (a, d_a, N * sizeof (float));
acc_map_data (b, d_b, N * sizeof (float));
exp = 3.0;
exp2 = 9.0;
#endif
return 0;
}