blob: 9ec345361d8e215b535d93e9e537fbd63e6f09df [file] [log] [blame]
/* acc_present_or_create, acc_present_or_copyin, etc. */
/* See also Fortran variants in "../libgomp.oacc-fortran/lib-32*". */
#include <stdbool.h>
#include <stdlib.h>
#include <openacc.h>
int
main (int argc, char **argv)
{
int *h, *d;
const int N = 10000;
const int S = N * sizeof *h;
bool shared_mem;
h = (int *) malloc (S);
if (!h)
abort ();
for (int i = 0; i < N; ++i)
h[i] = i + 0;
shared_mem = acc_is_present (h, S);
d = (int *) acc_present_or_create (h, S);
if (!d)
abort ();
if (shared_mem)
if (h != d)
abort ();
if (!acc_is_present (h, S))
abort ();
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
d[i] = i + 1;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 1 : 0))
abort ();
h[i] = i + 2;
}
{
int *d_ = (int *) acc_present_or_create (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 2 : 1))
abort ();
d[i] = i + 3;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 3 : 2))
abort ();
h[i] = i + 4;
}
{
int *d_ = (int *) acc_pcreate (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 4 : 3))
abort ();
d[i] = i + 5;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 5 : 4))
abort ();
h[i] = i + 6;
}
{
int *d_ = (int *) acc_present_or_copyin (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 6 : 5))
abort ();
d[i] = i + 7;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 7 : 6))
abort ();
h[i] = i + 8;
}
{
int *d_ = (int *) acc_pcopyin (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 8 : 7))
abort ();
d[i] = i + 9;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 9 : 8))
abort ();
h[i] = i + 10;
}
acc_copyout_finalize (h, S);
d = NULL;
if (!shared_mem)
if (acc_is_present (h, S))
abort ();
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 10 : 9))
abort ();
}
d = (int *) acc_pcopyin (h, S);
if (!d)
abort ();
if (shared_mem)
if (h != d)
abort ();
if (!acc_is_present (h, S))
abort ();
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 10 : 9))
abort ();
d[i] = i + 11;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 11 : 9))
abort ();
h[i] = i + 12;
}
{
int *d_ = (int *) acc_pcopyin (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 12 : 11))
abort ();
d[i] = i + 13;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 13 : 12))
abort ();
h[i] = i + 14;
}
{
int *d_ = (int *) acc_pcreate (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 14 : 13))
abort ();
d[i] = i + 15;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 15 : 14))
abort ();
h[i] = i + 16;
}
{
int *d_ = (int *) acc_pcreate (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 16 : 15))
abort ();
d[i] = i + 17;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 17 : 16))
abort ();
h[i] = i + 18;
}
acc_update_self (h, S);
if (!acc_is_present (h, S))
abort ();
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 18 : 17))
abort ();
}
acc_delete_finalize (h, S);
d = NULL;
if (!shared_mem)
if (acc_is_present (h, S))
abort();
free (h);
return 0;
}