int
main (int argc, char **argv)
{
const int N = 256;
int i;
unsigned char *h;
void *d;
h = (unsigned char *) malloc (N);
for (i = 0; i < N; i++)
{
h[i] = i;
}
d = acc_present_or_copyin (h, 0);
if (!d)
abort ();
memset (&h[0], 0, N);
acc_copyout (h, N);
for (i = 0; i < N; i++)
{
if (h[i] != i)
abort ();
}
free (h);
return 0;
}
int
main (int argc, char *argv[])
{
int i;
acc_present_or_copyin (&i, sizeof i);
fprintf (stderr, "CheCKpOInT\n");
#pragma acc enter data create (i)
return 0;
}
void
GOACC_enter_exit_data (int device, const void *offload_table, size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned short *kinds,
int async, int num_waits, ...)
{
struct goacc_thread *thr;
struct gomp_device_descr *acc_dev;
bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK;
bool data_enter = false;
size_t i;
select_acc_device (device);
thr = goacc_thread ();
acc_dev = thr->dev;
if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|| host_fallback)
return;
if (num_waits > 0)
{
va_list ap;
va_start (ap, num_waits);
goacc_wait (async, num_waits, ap);
va_end (ap);
}
acc_dev->openacc.async_set_async_func (async);
/* Determine if this is an "acc enter data". */
for (i = 0; i < mapnum; ++i)
{
unsigned char kind = kinds[i] & 0xff;
if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
continue;
if (kind == GOMP_MAP_FORCE_ALLOC
|| kind == GOMP_MAP_FORCE_PRESENT
|| kind == GOMP_MAP_FORCE_TO)
{
data_enter = true;
break;
}
if (kind == GOMP_MAP_FORCE_DEALLOC
|| kind == GOMP_MAP_FORCE_FROM)
break;
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
kind);
}
if (data_enter)
{
for (i = 0; i < mapnum; i++)
{
unsigned char kind = kinds[i] & 0xff;
/* Scan for PSETs. */
int psets = find_pset (i, mapnum, kinds);
if (!psets)
{
switch (kind)
{
case GOMP_MAP_POINTER:
gomp_acc_insert_pointer (1, &hostaddrs[i], &sizes[i],
&kinds[i]);
break;
case GOMP_MAP_FORCE_ALLOC:
acc_create (hostaddrs[i], sizes[i]);
break;
case GOMP_MAP_FORCE_PRESENT:
acc_present_or_copyin (hostaddrs[i], sizes[i]);
break;
case GOMP_MAP_FORCE_TO:
acc_present_or_copyin (hostaddrs[i], sizes[i]);
break;
default:
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
kind);
break;
}
}
else
{
gomp_acc_insert_pointer (3, &hostaddrs[i], &sizes[i], &kinds[i]);
/* Increment 'i' by two because OpenACC requires fortran
arrays to be contiguous, so each PSET is associated with
one of MAP_FORCE_ALLOC/MAP_FORCE_PRESET/MAP_FORCE_TO, and
one MAP_POINTER. */
i += 2;
}
}
}
else
for (i = 0; i < mapnum; ++i)
{
unsigned char kind = kinds[i] & 0xff;
int psets = find_pset (i, mapnum, kinds);
if (!psets)
{
switch (kind)
{
case GOMP_MAP_POINTER:
gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
== GOMP_MAP_FORCE_FROM,
async, 1);
break;
case GOMP_MAP_FORCE_DEALLOC:
acc_delete (hostaddrs[i], sizes[i]);
break;
case GOMP_MAP_FORCE_FROM:
acc_copyout (hostaddrs[i], sizes[i]);
break;
default:
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
kind);
break;
}
}
else
{
gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
== GOMP_MAP_FORCE_FROM, async, 3);
/* See the above comment. */
i += 2;
}
}
acc_dev->openacc.async_set_async_func (acc_async_sync);
}
